1//===- InferAddressSpace.cpp - --------------------------------------------===//
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// CUDA C/C++ includes memory space designation as variable type qualifers (such
10// as __global__ and __shared__). Knowing the space of a memory access allows
11// CUDA compilers to emit faster PTX loads and stores. For example, a load from
12// shared memory can be translated to `ld.shared` which is roughly 10% faster
13// than a generic `ld` on an NVIDIA Tesla K40c.
14//
15// Unfortunately, type qualifiers only apply to variable declarations, so CUDA
16// compilers must infer the memory space of an address expression from
17// type-qualified variables.
18//
19// LLVM IR uses non-zero (so-called) specific address spaces to represent memory
20// spaces (e.g. addrspace(3) means shared memory). The Clang frontend
21// places only type-qualified variables in specific address spaces, and then
22// conservatively `addrspacecast`s each type-qualified variable to addrspace(0)
23// (so-called the generic address space) for other instructions to use.
24//
25// For example, the Clang translates the following CUDA code
26// __shared__ float a[10];
27// float v = a[i];
28// to
29// %0 = addrspacecast [10 x float] addrspace(3)* @a to [10 x float]*
30// %1 = gep [10 x float], [10 x float]* %0, i64 0, i64 %i
31// %v = load float, float* %1 ; emits ld.f32
32// @a is in addrspace(3) since it's type-qualified, but its use from %1 is
33// redirected to %0 (the generic version of @a).
34//
35// The optimization implemented in this file propagates specific address spaces
36// from type-qualified variable declarations to its users. For example, it
37// optimizes the above IR to
38// %1 = gep [10 x float] addrspace(3)* @a, i64 0, i64 %i
39// %v = load float addrspace(3)* %1 ; emits ld.shared.f32
40// propagating the addrspace(3) from @a to %1. As the result, the NVPTX
41// codegen is able to emit ld.shared.f32 for %v.
42//
43// Address space inference works in two steps. First, it uses a data-flow
44// analysis to infer as many generic pointers as possible to point to only one
45// specific address space. In the above example, it can prove that %1 only
46// points to addrspace(3). This algorithm was published in
47// CUDA: Compiling and optimizing for a GPU platform
48// Chakrabarti, Grover, Aarts, Kong, Kudlur, Lin, Marathe, Murphy, Wang
49// ICCS 2012
50//
51// Then, address space inference replaces all refinable generic pointers with
52// equivalent specific pointers.
53//
54// The major challenge of implementing this optimization is handling PHINodes,
55// which may create loops in the data flow graph. This brings two complications.
56//
57// First, the data flow analysis in Step 1 needs to be circular. For example,
58// %generic.input = addrspacecast float addrspace(3)* %input to float*
59// loop:
60// %y = phi [ %generic.input, %y2 ]
61// %y2 = getelementptr %y, 1
62// %v = load %y2
63// br ..., label %loop, ...
64// proving %y specific requires proving both %generic.input and %y2 specific,
65// but proving %y2 specific circles back to %y. To address this complication,
66// the data flow analysis operates on a lattice:
67// uninitialized > specific address spaces > generic.
68// All address expressions (our implementation only considers phi, bitcast,
69// addrspacecast, and getelementptr) start with the uninitialized address space.
70// The monotone transfer function moves the address space of a pointer down a
71// lattice path from uninitialized to specific and then to generic. A join
72// operation of two different specific address spaces pushes the expression down
73// to their common address space. The analysis completes once it reaches a fixed
74// point.
75//
76// Second, IR rewriting in Step 2 also needs to be circular. For example,
77// converting %y to addrspace(3) requires the compiler to know the converted
78// %y2, but converting %y2 needs the converted %y. To address this complication,
79// we break these cycles using "poison" placeholders. When converting an
80// instruction `I` to a new address space, if its operand `Op` is not converted
81// yet, we let `I` temporarily use `poison` and fix all the uses later.
82// For instance, our algorithm first converts %y to
83// %y' = phi float addrspace(3)* [ %input, poison ]
84// Then, it converts %y2 to
85// %y2' = getelementptr %y', 1
86// Finally, it fixes the poison in %y' so that
87// %y' = phi float addrspace(3)* [ %input, %y2' ]
88//
89//===----------------------------------------------------------------------===//
90
91#include "llvm/Transforms/Scalar/InferAddressSpaces.h"
92#include "llvm/ADT/ArrayRef.h"
93#include "llvm/ADT/DenseMap.h"
94#include "llvm/ADT/DenseSet.h"
95#include "llvm/ADT/SetVector.h"
96#include "llvm/ADT/SmallVector.h"
97#include "llvm/Analysis/AssumptionCache.h"
98#include "llvm/Analysis/TargetTransformInfo.h"
99#include "llvm/Analysis/ValueTracking.h"
100#include "llvm/IR/Argument.h"
101#include "llvm/IR/BasicBlock.h"
102#include "llvm/IR/Constant.h"
103#include "llvm/IR/Constants.h"
104#include "llvm/IR/Dominators.h"
105#include "llvm/IR/Function.h"
106#include "llvm/IR/IRBuilder.h"
107#include "llvm/IR/InstIterator.h"
108#include "llvm/IR/Instruction.h"
109#include "llvm/IR/Instructions.h"
110#include "llvm/IR/IntrinsicInst.h"
111#include "llvm/IR/Intrinsics.h"
112#include "llvm/IR/LLVMContext.h"
113#include "llvm/IR/Operator.h"
114#include "llvm/IR/PassManager.h"
115#include "llvm/IR/PatternMatch.h"
116#include "llvm/IR/Type.h"
117#include "llvm/IR/Use.h"
118#include "llvm/IR/User.h"
119#include "llvm/IR/Value.h"
120#include "llvm/IR/ValueHandle.h"
121#include "llvm/InitializePasses.h"
122#include "llvm/Pass.h"
123#include "llvm/Support/Casting.h"
124#include "llvm/Support/Debug.h"
125#include "llvm/Support/ErrorHandling.h"
126#include "llvm/Support/KnownBits.h"
127#include "llvm/Support/raw_ostream.h"
128#include "llvm/Transforms/Scalar.h"
129#include "llvm/Transforms/Utils/Local.h"
130#include "llvm/Transforms/Utils/ValueMapper.h"
131#include <cassert>
132#include <iterator>
133#include <limits>
134#include <optional>
135#include <utility>
136#include <vector>
137
138#define DEBUG_TYPE "infer-address-spaces"
139
140using namespace llvm;
141using namespace llvm::PatternMatch;
142
143static const unsigned UninitializedAddressSpace =
144 std::numeric_limits<unsigned>::max();
145
146namespace {
147
148using ValueToAddrSpaceMapTy = DenseMap<const Value *, unsigned>;
149// Different from ValueToAddrSpaceMapTy, where a new addrspace is inferred on
150// the *def* of a value, PredicatedAddrSpaceMapTy is map where a new
151// addrspace is inferred on the *use* of a pointer. This map is introduced to
152// infer addrspace from the addrspace predicate assumption built from assume
153// intrinsic. In that scenario, only specific uses (under valid assumption
154// context) could be inferred with a new addrspace.
155using PredicatedAddrSpaceMapTy =
156 DenseMap<std::pair<const Value *, const Value *>, unsigned>;
157using PostorderStackTy = llvm::SmallVector<PointerIntPair<Value *, 1, bool>, 4>;
158
159class InferAddressSpaces : public FunctionPass {
160 unsigned FlatAddrSpace = 0;
161
162public:
163 static char ID;
164
165 InferAddressSpaces()
166 : FunctionPass(ID), FlatAddrSpace(UninitializedAddressSpace) {
167 initializeInferAddressSpacesPass(*PassRegistry::getPassRegistry());
168 }
169 InferAddressSpaces(unsigned AS) : FunctionPass(ID), FlatAddrSpace(AS) {
170 initializeInferAddressSpacesPass(*PassRegistry::getPassRegistry());
171 }
172
173 void getAnalysisUsage(AnalysisUsage &AU) const override {
174 AU.setPreservesCFG();
175 AU.addRequired<AssumptionCacheTracker>();
176 AU.addRequired<TargetTransformInfoWrapperPass>();
177 }
178
179 bool runOnFunction(Function &F) override;
180};
181
182class InferAddressSpacesImpl {
183 AssumptionCache &AC;
184 Function *F = nullptr;
185 const DominatorTree *DT = nullptr;
186 const TargetTransformInfo *TTI = nullptr;
187 const DataLayout *DL = nullptr;
188
189 /// Target specific address space which uses of should be replaced if
190 /// possible.
191 unsigned FlatAddrSpace = 0;
192
193 /// The default address space is assumed as the flat address space. This is
194 /// mainly for test purpose.
195 const bool AssumeDefaultIsFlatAddressSpace = false;
196
197 DenseMap<const Value *, Value *> PtrIntCastPairs;
198
199 // Tries to find if the inttoptr instruction is derived from an pointer have
200 // specific address space, and is safe to propagate the address space to the
201 // new pointer that inttoptr produces.
202 Value *getIntToPtrPointerOperand(const Operator *I2P) const;
203 // Tries to find if the inttoptr instruction is derived from an pointer have
204 // specific address space, and is safe to propagate the address space to the
205 // new pointer that inttoptr produces. If the old pointer is found, cache the
206 // <OldPtr, inttoptr> pairs to a map.
207 void collectIntToPtrPointerOperand();
208 // Check if an old pointer is found ahead of time. The safety has been checked
209 // when collecting the inttoptr original pointer and the result is cached in
210 // PtrIntCastPairs.
211 bool isSafeToCastIntToPtrAddrSpace(const Operator *I2P) const {
212 return PtrIntCastPairs.contains(Val: I2P);
213 }
214 bool isAddressExpression(const Value &V, const DataLayout &DL,
215 const TargetTransformInfo *TTI) const;
216 Value *cloneConstantExprWithNewAddressSpace(
217 ConstantExpr *CE, unsigned NewAddrSpace,
218 const ValueToValueMapTy &ValueWithNewAddrSpace, const DataLayout *DL,
219 const TargetTransformInfo *TTI) const;
220
221 SmallVector<Value *, 2>
222 getPointerOperands(const Value &V, const DataLayout &DL,
223 const TargetTransformInfo *TTI) const;
224
225 // Try to update the address space of V. If V is updated, returns true and
226 // false otherwise.
227 bool updateAddressSpace(const Value &V,
228 ValueToAddrSpaceMapTy &InferredAddrSpace,
229 PredicatedAddrSpaceMapTy &PredicatedAS) const;
230
231 // Adds the users of V whose address space may still change to Worklist.
232 void enqueueUsers(Value &V, const ValueToAddrSpaceMapTy &InferredAddrSpace,
233 SetVector<Value *> &Worklist) const;
234
235 // Propagates address spaces out of Worklist until nothing changes.
236 void runToFixPoint(SetVector<Value *> &Worklist,
237 ValueToAddrSpaceMapTy &InferredAddrSpace,
238 PredicatedAddrSpaceMapTy &PredicatedAS) const;
239
240 // Tries to infer the specific address space of each address expression in
241 // Postorder.
242 void inferAddressSpaces(ArrayRef<WeakTrackingVH> Postorder,
243 ValueToAddrSpaceMapTy &InferredAddrSpace,
244 PredicatedAddrSpaceMapTy &PredicatedAS) const;
245
246 bool isSafeToCastConstAddrSpace(Constant *C, unsigned NewAS) const;
247
248 Value *clonePtrMaskWithNewAddressSpace(
249 IntrinsicInst *I, unsigned NewAddrSpace,
250 const ValueToValueMapTy &ValueWithNewAddrSpace,
251 const PredicatedAddrSpaceMapTy &PredicatedAS,
252 SmallVectorImpl<const Use *> *PoisonUsesToFix) const;
253
254 Value *cloneInstructionWithNewAddressSpace(
255 Instruction *I, unsigned NewAddrSpace,
256 const ValueToValueMapTy &ValueWithNewAddrSpace,
257 const PredicatedAddrSpaceMapTy &PredicatedAS,
258 SmallVectorImpl<const Use *> *PoisonUsesToFix) const;
259
260 void performPointerReplacement(
261 Value *V, Value *NewV, Use &U, ValueToValueMapTy &ValueWithNewAddrSpace,
262 SmallVectorImpl<Instruction *> &DeadInstructions) const;
263
264 // Changes the flat address expressions in function F to point to specific
265 // address spaces if InferredAddrSpace says so. Postorder is the postorder of
266 // all flat expressions in the use-def graph of function F.
267 bool rewriteWithNewAddressSpaces(
268 ArrayRef<WeakTrackingVH> Postorder,
269 const ValueToAddrSpaceMapTy &InferredAddrSpace,
270 const PredicatedAddrSpaceMapTy &PredicatedAS) const;
271
272 void appendsFlatAddressExpressionToPostorderStack(
273 Value *V, PostorderStackTy &PostorderStack,
274 DenseSet<Value *> &Visited) const;
275
276 bool rewriteIntrinsicOperands(IntrinsicInst *II, Value *OldV,
277 Value *NewV) const;
278 void collectRewritableIntrinsicOperands(IntrinsicInst *II,
279 PostorderStackTy &PostorderStack,
280 DenseSet<Value *> &Visited) const;
281
282 std::vector<WeakTrackingVH> collectFlatAddressExpressions(Function &F) const;
283
284 Value *cloneValueWithNewAddressSpace(
285 Value *V, unsigned NewAddrSpace,
286 const ValueToValueMapTy &ValueWithNewAddrSpace,
287 const PredicatedAddrSpaceMapTy &PredicatedAS,
288 SmallVectorImpl<const Use *> *PoisonUsesToFix) const;
289 unsigned joinAddressSpaces(unsigned AS1, unsigned AS2) const;
290
291 unsigned getPredicatedAddrSpace(const Value &PtrV,
292 const Value *UserCtx) const;
293
294public:
295 InferAddressSpacesImpl(AssumptionCache &AC, const DominatorTree *DT,
296 const TargetTransformInfo *TTI, unsigned FlatAddrSpace,
297 bool AssumeDefaultIsFlatAddressSpace)
298 : AC(AC), DT(DT), TTI(TTI), FlatAddrSpace(FlatAddrSpace),
299 AssumeDefaultIsFlatAddressSpace(AssumeDefaultIsFlatAddressSpace) {}
300 bool run(Function &F);
301};
302
303} // end anonymous namespace
304
305char InferAddressSpaces::ID = 0;
306
307INITIALIZE_PASS_BEGIN(InferAddressSpaces, DEBUG_TYPE, "Infer address spaces",
308 false, false)
309INITIALIZE_PASS_DEPENDENCY(AssumptionCacheTracker)
310INITIALIZE_PASS_DEPENDENCY(TargetTransformInfoWrapperPass)
311INITIALIZE_PASS_END(InferAddressSpaces, DEBUG_TYPE, "Infer address spaces",
312 false, false)
313
314static Type *getPtrOrVecOfPtrsWithNewAS(Type *Ty, unsigned NewAddrSpace) {
315 assert(Ty->isPtrOrPtrVectorTy());
316 PointerType *NPT = PointerType::get(C&: Ty->getContext(), AddressSpace: NewAddrSpace);
317 return Ty->getWithNewType(EltTy: NPT);
318}
319
320// Check whether that's no-op pointer bitcast using a pair of
321// `ptrtoint`/`inttoptr` due to the missing no-op pointer bitcast over
322// different address spaces.
323static bool isNoopPtrIntCastPair(const Operator *I2P, const DataLayout &DL,
324 const TargetTransformInfo *TTI) {
325 assert(I2P->getOpcode() == Instruction::IntToPtr);
326 auto *P2I = dyn_cast<Operator>(Val: I2P->getOperand(i: 0));
327 if (!P2I || P2I->getOpcode() != Instruction::PtrToInt)
328 return false;
329 // Check it's really safe to treat that pair of `ptrtoint`/`inttoptr` as a
330 // no-op cast. Besides checking both of them are no-op casts, as the
331 // reinterpreted pointer may be used in other pointer arithmetic, we also
332 // need to double-check that through the target-specific hook. That ensures
333 // the underlying target also agrees that's a no-op address space cast and
334 // pointer bits are preserved.
335 // The current IR spec doesn't have clear rules on address space casts,
336 // especially a clear definition for pointer bits in non-default address
337 // spaces. It would be undefined if that pointer is dereferenced after an
338 // invalid reinterpret cast. Also, due to the unclearness for the meaning of
339 // bits in non-default address spaces in the current spec, the pointer
340 // arithmetic may also be undefined after invalid pointer reinterpret cast.
341 // However, as we confirm through the target hooks that it's a no-op
342 // addrspacecast, it doesn't matter since the bits should be the same.
343 unsigned P2IOp0AS = P2I->getOperand(i: 0)->getType()->getPointerAddressSpace();
344 unsigned I2PAS = I2P->getType()->getPointerAddressSpace();
345 return CastInst::isNoopCast(Opcode: Instruction::CastOps(I2P->getOpcode()),
346 SrcTy: I2P->getOperand(i: 0)->getType(), DstTy: I2P->getType(),
347 DL) &&
348 CastInst::isNoopCast(Opcode: Instruction::CastOps(P2I->getOpcode()),
349 SrcTy: P2I->getOperand(i: 0)->getType(), DstTy: P2I->getType(),
350 DL) &&
351 (P2IOp0AS == I2PAS || TTI->isNoopAddrSpaceCast(FromAS: P2IOp0AS, ToAS: I2PAS));
352}
353
354// Returns true if V is an address expression.
355// TODO: Currently, we only consider:
356// - arguments
357// - phi, bitcast, addrspacecast, and getelementptr operators
358bool InferAddressSpacesImpl::isAddressExpression(
359 const Value &V, const DataLayout &DL,
360 const TargetTransformInfo *TTI) const {
361
362 if (const Argument *Arg = dyn_cast<Argument>(Val: &V))
363 return Arg->getType()->isPointerTy() &&
364 TTI->getAssumedAddrSpace(V: &V) != UninitializedAddressSpace;
365
366 const Operator *Op = dyn_cast<Operator>(Val: &V);
367 if (!Op)
368 return false;
369
370 switch (Op->getOpcode()) {
371 case Instruction::PHI:
372 assert(Op->getType()->isPtrOrPtrVectorTy());
373 return true;
374 case Instruction::BitCast:
375 case Instruction::AddrSpaceCast:
376 case Instruction::GetElementPtr:
377 return true;
378 case Instruction::Select:
379 return Op->getType()->isPtrOrPtrVectorTy();
380 case Instruction::Call: {
381 const IntrinsicInst *II = dyn_cast<IntrinsicInst>(Val: &V);
382 return II && II->getIntrinsicID() == Intrinsic::ptrmask;
383 }
384 case Instruction::IntToPtr:
385 return isNoopPtrIntCastPair(I2P: Op, DL, TTI) ||
386 isSafeToCastIntToPtrAddrSpace(I2P: Op);
387 default:
388 // That value is an address expression if it has an assumed address space.
389 return TTI->getAssumedAddrSpace(V: &V) != UninitializedAddressSpace;
390 }
391}
392
393// Returns the pointer operands of V.
394//
395// Precondition: V is an address expression.
396SmallVector<Value *, 2> InferAddressSpacesImpl::getPointerOperands(
397 const Value &V, const DataLayout &DL,
398 const TargetTransformInfo *TTI) const {
399 if (isa<Argument>(Val: &V))
400 return {};
401
402 const Operator &Op = cast<Operator>(Val: V);
403 switch (Op.getOpcode()) {
404 case Instruction::PHI: {
405 auto IncomingValues = cast<PHINode>(Val: Op).incoming_values();
406 return {IncomingValues.begin(), IncomingValues.end()};
407 }
408 case Instruction::BitCast:
409 case Instruction::AddrSpaceCast:
410 case Instruction::GetElementPtr:
411 return {Op.getOperand(i: 0)};
412 case Instruction::Select:
413 return {Op.getOperand(i: 1), Op.getOperand(i: 2)};
414 case Instruction::Call: {
415 const IntrinsicInst &II = cast<IntrinsicInst>(Val: Op);
416 assert(II.getIntrinsicID() == Intrinsic::ptrmask &&
417 "unexpected intrinsic call");
418 return {II.getArgOperand(i: 0)};
419 }
420 case Instruction::IntToPtr: {
421 if (isNoopPtrIntCastPair(I2P: &Op, DL, TTI)) {
422 auto *P2I = cast<Operator>(Val: Op.getOperand(i: 0));
423 return {P2I->getOperand(i: 0)};
424 }
425 assert(isSafeToCastIntToPtrAddrSpace(&Op));
426 return {getIntToPtrPointerOperand(I2P: &Op)};
427 }
428 default:
429 llvm_unreachable("Unexpected instruction type.");
430 }
431}
432
433// Return mask. The 1 in mask indicate the bit is changed.
434// This helper function is to compute the max know changed bits for ptr1 and
435// ptr2 after the operation `ptr2 = ptr1 Op Mask`.
436static APInt computeMaxChangedPtrBits(const Operator *Op, const Value *Mask,
437 const DataLayout &DL, AssumptionCache *AC,
438 const DominatorTree *DT) {
439 KnownBits Known = computeKnownBits(V: Mask, DL, AC, CtxI: nullptr, DT);
440 switch (Op->getOpcode()) {
441 case Instruction::Xor:
442 case Instruction::Or:
443 return ~Known.Zero;
444 case Instruction::And:
445 return ~Known.One;
446 default:
447 return APInt::getAllOnes(numBits: Known.getBitWidth());
448 }
449}
450
451Value *
452InferAddressSpacesImpl::getIntToPtrPointerOperand(const Operator *I2P) const {
453 assert(I2P->getOpcode() == Instruction::IntToPtr);
454 if (I2P->getType()->isVectorTy())
455 return nullptr;
456
457 // If I2P has been accessed and has the corresponding old pointer value, just
458 // return true.
459 if (auto *OldPtr = PtrIntCastPairs.lookup(Val: I2P))
460 return OldPtr;
461
462 Value *LogicalOp = I2P->getOperand(i: 0);
463 Value *OldPtr, *Mask;
464 if (!match(V: LogicalOp,
465 P: m_c_BitwiseLogic(L: m_PtrToInt(Op: m_Value(V&: OldPtr)), R: m_Value(V&: Mask))))
466 return nullptr;
467
468 Operator *AsCast = dyn_cast<AddrSpaceCastOperator>(Val: OldPtr);
469 if (!AsCast)
470 return nullptr;
471
472 unsigned SrcAS = I2P->getType()->getPointerAddressSpace();
473 unsigned DstAS = AsCast->getOperand(i: 0)->getType()->getPointerAddressSpace();
474 APInt PreservedPtrMask = TTI->getAddrSpaceCastPreservedPtrMask(SrcAS, DstAS);
475 if (PreservedPtrMask.isZero())
476 return nullptr;
477 APInt ChangedPtrBits =
478 computeMaxChangedPtrBits(Op: cast<Operator>(Val: LogicalOp), Mask, DL: *DL, AC: &AC, DT);
479 // Check if the address bits change is within the preserved mask. If the bits
480 // change is not preserved, it is not safe to perform address space cast.
481 // The following pattern is not safe to cast address space.
482 // %1 = ptrtoint ptr addrspace(3) %sp to i32
483 // %2 = zext i32 %1 to i64
484 // %gp = inttoptr i64 %2 to ptr
485 assert(ChangedPtrBits.getBitWidth() == PreservedPtrMask.getBitWidth());
486 if (ChangedPtrBits.isSubsetOf(RHS: PreservedPtrMask))
487 return OldPtr;
488
489 return nullptr;
490}
491
492void InferAddressSpacesImpl::collectIntToPtrPointerOperand() {
493 // Only collect inttoptr instruction.
494 // TODO: We need to collect inttoptr constant expression as well.
495 for (Instruction &I : instructions(F)) {
496 if (!dyn_cast<IntToPtrInst>(Val: &I))
497 continue;
498 if (auto *OldPtr = getIntToPtrPointerOperand(I2P: cast<Operator>(Val: &I)))
499 PtrIntCastPairs.insert(KV: {&I, OldPtr});
500 }
501}
502
503bool InferAddressSpacesImpl::rewriteIntrinsicOperands(IntrinsicInst *II,
504 Value *OldV,
505 Value *NewV) const {
506 Module *M = II->getParent()->getParent()->getParent();
507 Intrinsic::ID IID = II->getIntrinsicID();
508 switch (IID) {
509 case Intrinsic::objectsize:
510 case Intrinsic::masked_load: {
511 Type *DestTy = II->getType();
512 Type *SrcTy = NewV->getType();
513 Function *NewDecl =
514 Intrinsic::getOrInsertDeclaration(M, id: IID, OverloadTys: {DestTy, SrcTy});
515 II->setArgOperand(i: 0, v: NewV);
516 II->setCalledFunction(NewDecl);
517 return true;
518 }
519 case Intrinsic::ptrmask:
520 // This is handled as an address expression, not as a use memory operation.
521 return false;
522 case Intrinsic::masked_gather: {
523 Type *RetTy = II->getType();
524 Type *NewPtrTy = NewV->getType();
525 Function *NewDecl =
526 Intrinsic::getOrInsertDeclaration(M, id: IID, OverloadTys: {RetTy, NewPtrTy});
527 II->setArgOperand(i: 0, v: NewV);
528 II->setCalledFunction(NewDecl);
529 return true;
530 }
531 case Intrinsic::masked_store:
532 case Intrinsic::masked_scatter: {
533 Type *ValueTy = II->getOperand(i_nocapture: 0)->getType();
534 Type *NewPtrTy = NewV->getType();
535 Function *NewDecl = Intrinsic::getOrInsertDeclaration(
536 M, id: II->getIntrinsicID(), OverloadTys: {ValueTy, NewPtrTy});
537 II->setArgOperand(i: 1, v: NewV);
538 II->setCalledFunction(NewDecl);
539 return true;
540 }
541 case Intrinsic::prefetch:
542 case Intrinsic::is_constant: {
543 Function *NewDecl = Intrinsic::getOrInsertDeclaration(
544 M, id: II->getIntrinsicID(), OverloadTys: {NewV->getType()});
545 II->setArgOperand(i: 0, v: NewV);
546 II->setCalledFunction(NewDecl);
547 return true;
548 }
549 case Intrinsic::fake_use: {
550 II->replaceUsesOfWith(From: OldV, To: NewV);
551 return true;
552 }
553 case Intrinsic::lifetime_start:
554 case Intrinsic::lifetime_end: {
555 // Always force lifetime markers to work directly on the alloca.
556 NewV = NewV->stripPointerCasts();
557 Function *NewDecl = Intrinsic::getOrInsertDeclaration(
558 M, id: II->getIntrinsicID(), OverloadTys: {NewV->getType()});
559 II->setArgOperand(i: 0, v: NewV);
560 II->setCalledFunction(NewDecl);
561 return true;
562 }
563 default: {
564 Value *Rewrite = TTI->rewriteIntrinsicWithAddressSpace(II, OldV, NewV);
565 if (!Rewrite)
566 return false;
567 if (Rewrite != II)
568 II->replaceAllUsesWith(V: Rewrite);
569 return true;
570 }
571 }
572}
573
574void InferAddressSpacesImpl::collectRewritableIntrinsicOperands(
575 IntrinsicInst *II, PostorderStackTy &PostorderStack,
576 DenseSet<Value *> &Visited) const {
577 auto IID = II->getIntrinsicID();
578 switch (IID) {
579 case Intrinsic::ptrmask:
580 case Intrinsic::objectsize:
581 appendsFlatAddressExpressionToPostorderStack(V: II->getArgOperand(i: 0),
582 PostorderStack, Visited);
583 break;
584 case Intrinsic::is_constant: {
585 Value *Ptr = II->getArgOperand(i: 0);
586 if (Ptr->getType()->isPtrOrPtrVectorTy()) {
587 appendsFlatAddressExpressionToPostorderStack(V: Ptr, PostorderStack,
588 Visited);
589 }
590
591 break;
592 }
593 case Intrinsic::masked_load:
594 case Intrinsic::masked_gather:
595 case Intrinsic::prefetch:
596 appendsFlatAddressExpressionToPostorderStack(V: II->getArgOperand(i: 0),
597 PostorderStack, Visited);
598 break;
599 case Intrinsic::masked_store:
600 case Intrinsic::masked_scatter:
601 appendsFlatAddressExpressionToPostorderStack(V: II->getArgOperand(i: 1),
602 PostorderStack, Visited);
603 break;
604 case Intrinsic::fake_use: {
605 for (Value *Op : II->operands()) {
606 if (Op->getType()->isPtrOrPtrVectorTy()) {
607 appendsFlatAddressExpressionToPostorderStack(V: Op, PostorderStack,
608 Visited);
609 }
610 }
611
612 break;
613 }
614 case Intrinsic::lifetime_start:
615 case Intrinsic::lifetime_end: {
616 appendsFlatAddressExpressionToPostorderStack(V: II->getArgOperand(i: 0),
617 PostorderStack, Visited);
618 break;
619 }
620 default:
621 SmallVector<int, 2> OpIndexes;
622 if (TTI->collectFlatAddressOperands(OpIndexes, IID)) {
623 for (int Idx : OpIndexes) {
624 appendsFlatAddressExpressionToPostorderStack(V: II->getArgOperand(i: Idx),
625 PostorderStack, Visited);
626 }
627 }
628 break;
629 }
630}
631
632// Returns all flat address expressions in function F. The elements are
633// If V is an unvisited flat address expression, appends V to PostorderStack
634// and marks it as visited.
635void InferAddressSpacesImpl::appendsFlatAddressExpressionToPostorderStack(
636 Value *V, PostorderStackTy &PostorderStack,
637 DenseSet<Value *> &Visited) const {
638 assert(V->getType()->isPtrOrPtrVectorTy());
639
640 // Generic addressing expressions may be hidden in nested constant
641 // expressions.
642 if (ConstantExpr *CE = dyn_cast<ConstantExpr>(Val: V)) {
643 // TODO: Look in non-address parts, like icmp operands.
644 if (isAddressExpression(V: *CE, DL: *DL, TTI) && Visited.insert(V: CE).second)
645 PostorderStack.emplace_back(Args&: CE, Args: false);
646
647 return;
648 }
649
650 if (V->getType()->getPointerAddressSpace() == FlatAddrSpace &&
651 isAddressExpression(V: *V, DL: *DL, TTI)) {
652 if (Visited.insert(V).second) {
653 PostorderStack.emplace_back(Args&: V, Args: false);
654
655 if (auto *Op = dyn_cast<Operator>(Val: V))
656 for (auto &O : Op->operands())
657 if (ConstantExpr *CE = dyn_cast<ConstantExpr>(Val&: O))
658 if (isAddressExpression(V: *CE, DL: *DL, TTI) && Visited.insert(V: CE).second)
659 PostorderStack.emplace_back(Args&: CE, Args: false);
660 }
661 }
662}
663
664// Returns all flat address expressions in function F. The elements are ordered
665// in postorder.
666std::vector<WeakTrackingVH>
667InferAddressSpacesImpl::collectFlatAddressExpressions(Function &F) const {
668 // This function implements a non-recursive postorder traversal of a partial
669 // use-def graph of function F.
670 PostorderStackTy PostorderStack;
671 // The set of visited expressions.
672 DenseSet<Value *> Visited;
673
674 auto PushPtrOperand = [&](Value *Ptr) {
675 appendsFlatAddressExpressionToPostorderStack(V: Ptr, PostorderStack, Visited);
676 };
677
678 // Look at operations that may be interesting accelerate by moving to a known
679 // address space. We aim at generating after loads and stores, but pure
680 // addressing calculations may also be faster.
681 for (Instruction &I : instructions(F)) {
682 if (auto *GEP = dyn_cast<GetElementPtrInst>(Val: &I)) {
683 PushPtrOperand(GEP->getPointerOperand());
684 } else if (auto *LI = dyn_cast<LoadInst>(Val: &I))
685 PushPtrOperand(LI->getPointerOperand());
686 else if (auto *SI = dyn_cast<StoreInst>(Val: &I))
687 PushPtrOperand(SI->getPointerOperand());
688 else if (auto *RMW = dyn_cast<AtomicRMWInst>(Val: &I))
689 PushPtrOperand(RMW->getPointerOperand());
690 else if (auto *CmpX = dyn_cast<AtomicCmpXchgInst>(Val: &I))
691 PushPtrOperand(CmpX->getPointerOperand());
692 else if (auto *MI = dyn_cast<MemIntrinsic>(Val: &I)) {
693 // For memset/memcpy/memmove, any pointer operand can be replaced.
694 PushPtrOperand(MI->getRawDest());
695
696 // Handle 2nd operand for memcpy/memmove.
697 if (auto *MTI = dyn_cast<MemTransferInst>(Val: MI))
698 PushPtrOperand(MTI->getRawSource());
699 } else if (auto *II = dyn_cast<IntrinsicInst>(Val: &I))
700 collectRewritableIntrinsicOperands(II, PostorderStack, Visited);
701 else if (ICmpInst *Cmp = dyn_cast<ICmpInst>(Val: &I)) {
702 if (Cmp->getOperand(i_nocapture: 0)->getType()->isPtrOrPtrVectorTy()) {
703 PushPtrOperand(Cmp->getOperand(i_nocapture: 0));
704 PushPtrOperand(Cmp->getOperand(i_nocapture: 1));
705 }
706 } else if (auto *ASC = dyn_cast<AddrSpaceCastInst>(Val: &I)) {
707 PushPtrOperand(ASC->getPointerOperand());
708 } else if (auto *I2P = dyn_cast<IntToPtrInst>(Val: &I)) {
709 if (isNoopPtrIntCastPair(I2P: cast<Operator>(Val: I2P), DL: *DL, TTI))
710 PushPtrOperand(cast<Operator>(Val: I2P->getOperand(i_nocapture: 0))->getOperand(i: 0));
711 else if (isSafeToCastIntToPtrAddrSpace(I2P: cast<Operator>(Val: I2P)))
712 PushPtrOperand(getIntToPtrPointerOperand(I2P: cast<Operator>(Val: I2P)));
713 } else if (auto *RI = dyn_cast<ReturnInst>(Val: &I)) {
714 if (auto *RV = RI->getReturnValue();
715 RV && RV->getType()->isPtrOrPtrVectorTy())
716 PushPtrOperand(RV);
717 }
718 }
719
720 std::vector<WeakTrackingVH> Postorder; // The resultant postorder.
721 while (!PostorderStack.empty()) {
722 Value *TopVal = PostorderStack.back().getPointer();
723 // If the operands of the expression on the top are already explored,
724 // adds that expression to the resultant postorder.
725 if (PostorderStack.back().getInt()) {
726 if (TopVal->getType()->getPointerAddressSpace() == FlatAddrSpace)
727 Postorder.push_back(x: TopVal);
728 PostorderStack.pop_back();
729 continue;
730 }
731 // Otherwise, adds its operands to the stack and explores them.
732 PostorderStack.back().setInt(true);
733 // Skip values with an assumed address space.
734 if (TTI->getAssumedAddrSpace(V: TopVal) == UninitializedAddressSpace) {
735 for (Value *PtrOperand : getPointerOperands(V: *TopVal, DL: *DL, TTI)) {
736 appendsFlatAddressExpressionToPostorderStack(V: PtrOperand, PostorderStack,
737 Visited);
738 }
739 }
740 }
741 return Postorder;
742}
743
744// Inserts an addrspacecast for a phi node operand, handling the proper
745// insertion position based on the operand type.
746static Value *phiNodeOperandWithNewAddressSpace(AddrSpaceCastInst *NewI,
747 Value *Operand) {
748 auto InsertBefore = [NewI](auto It) {
749 NewI->insertBefore(It);
750 NewI->setDebugLoc(It->getDebugLoc());
751 return NewI;
752 };
753
754 if (auto *Arg = dyn_cast<Argument>(Val: Operand)) {
755 // For arguments, insert the cast at the beginning of entry block.
756 // Consider inserting at the dominating block for better placement.
757 Function *F = Arg->getParent();
758 auto InsertI = F->getEntryBlock().getFirstNonPHIIt();
759 return InsertBefore(InsertI);
760 }
761
762 // No check for Constant here, as constants are already handled.
763 assert(isa<Instruction>(Operand));
764
765 Instruction *OpInst = cast<Instruction>(Val: Operand);
766 if (LLVM_UNLIKELY(OpInst->getOpcode() == Instruction::PHI)) {
767 // If the operand is defined by another PHI node, insert after the first
768 // non-PHI instruction at the corresponding basic block.
769 auto InsertI = OpInst->getParent()->getFirstNonPHIIt();
770 return InsertBefore(InsertI);
771 }
772
773 // Otherwise, insert immediately after the operand definition.
774 NewI->insertAfter(InsertPos: OpInst->getIterator());
775 NewI->setDebugLoc(OpInst->getDebugLoc());
776 return NewI;
777}
778
779// A helper function for cloneInstructionWithNewAddressSpace. Returns the clone
780// of OperandUse.get() in the new address space. If the clone is not ready yet,
781// returns poison in the new address space as a placeholder.
782static Value *operandWithNewAddressSpaceOrCreatePoison(
783 const Use &OperandUse, unsigned NewAddrSpace,
784 const ValueToValueMapTy &ValueWithNewAddrSpace,
785 const PredicatedAddrSpaceMapTy &PredicatedAS,
786 SmallVectorImpl<const Use *> *PoisonUsesToFix) {
787 Value *Operand = OperandUse.get();
788
789 Type *NewPtrTy = getPtrOrVecOfPtrsWithNewAS(Ty: Operand->getType(), NewAddrSpace);
790
791 if (Constant *C = dyn_cast<Constant>(Val: Operand))
792 return ConstantExpr::getAddrSpaceCast(C, Ty: NewPtrTy);
793
794 Instruction *Inst = cast<Instruction>(Val: OperandUse.getUser());
795 if (Value *NewOperand = ValueWithNewAddrSpace.lookup(Val: Operand)) {
796 Operand = NewOperand;
797 } else if (!PredicatedAS.contains(Val: std::make_pair(x&: Inst, y&: Operand))) {
798 assert(PoisonUsesToFix && "missing inferred operand replacement");
799 PoisonUsesToFix->push_back(Elt: &OperandUse);
800 return PoisonValue::get(T: NewPtrTy);
801 }
802
803 if (Operand->getType() == NewPtrTy)
804 return Operand;
805
806 auto *NewI = new AddrSpaceCastInst(Operand, NewPtrTy);
807 if (LLVM_UNLIKELY(Inst->getOpcode() == Instruction::PHI))
808 return phiNodeOperandWithNewAddressSpace(NewI, Operand: OperandUse.get());
809
810 // During cloning phase, the cast is placed before the original flat
811 // instruction, as its clone in the new address space has not been inserted
812 // yet. During poison fixup phase, the clone already exists, thus make sure
813 // the cast is inserted before it.
814 Instruction *InsertPt = Inst;
815 if (Value *NewUser = ValueWithNewAddrSpace.lookup(Val: Inst))
816 InsertPt = cast<Instruction>(Val: NewUser);
817 NewI->insertBefore(InsertPos: InsertPt->getIterator());
818 NewI->setDebugLoc(Inst->getDebugLoc());
819 return NewI;
820}
821
822// A helper function for cloneInstructionWithNewAddressSpace. Handles the
823// conversion of a ptrmask intrinsic instruction.
824Value *InferAddressSpacesImpl::clonePtrMaskWithNewAddressSpace(
825 IntrinsicInst *I, unsigned NewAddrSpace,
826 const ValueToValueMapTy &ValueWithNewAddrSpace,
827 const PredicatedAddrSpaceMapTy &PredicatedAS,
828 SmallVectorImpl<const Use *> *PoisonUsesToFix) const {
829 const Use &PtrOpUse = I->getArgOperandUse(i: 0);
830 unsigned OldAddrSpace = PtrOpUse->getType()->getPointerAddressSpace();
831 Value *MaskOp = I->getArgOperand(i: 1);
832 Type *MaskTy = MaskOp->getType();
833
834 KnownBits OldPtrBits{DL->getPointerSizeInBits(AS: OldAddrSpace)};
835 KnownBits NewPtrBits{DL->getPointerSizeInBits(AS: NewAddrSpace)};
836 if (!TTI->isNoopAddrSpaceCast(FromAS: OldAddrSpace, ToAS: NewAddrSpace)) {
837 std::tie(args&: OldPtrBits, args&: NewPtrBits) =
838 TTI->computeKnownBitsAddrSpaceCast(ToAS: NewAddrSpace, PtrOp: *PtrOpUse.get());
839 }
840
841 // If the pointers in both addrspaces have a bitwise representation and if the
842 // representation of the new pointer is smaller (fewer bits) than the old one,
843 // check if the mask is applicable to the ptr in the new addrspace. Any
844 // masking only clearing the low bits will also apply in the new addrspace
845 // Note: checking if the mask clears high bits is not sufficient as those
846 // might have already been 0 in the old ptr.
847 if (OldPtrBits.getBitWidth() > NewPtrBits.getBitWidth()) {
848 KnownBits MaskBits =
849 computeKnownBits(V: MaskOp, DL: *DL, /*AssumptionCache=*/AC: nullptr, CtxI: I);
850 // Set all unknown bits of the old ptr to 1, so that we are conservative in
851 // checking which bits are cleared by the mask.
852 OldPtrBits.One |= ~OldPtrBits.Zero;
853 // Check which bits are cleared by the mask in the old ptr.
854 KnownBits ClearedBits = KnownBits::sub(LHS: OldPtrBits, RHS: OldPtrBits & MaskBits);
855
856 // If the mask isn't applicable to the new ptr, leave the ptrmask as-is and
857 // insert an addrspacecast after it.
858 if (ClearedBits.countMaxActiveBits() > NewPtrBits.countMaxActiveBits()) {
859 std::optional<BasicBlock::iterator> InsertPoint =
860 I->getInsertionPointAfterDef();
861 assert(InsertPoint && "insertion after ptrmask should be possible");
862 Type *NewPtrType = getPtrOrVecOfPtrsWithNewAS(Ty: I->getType(), NewAddrSpace);
863 Instruction *AddrSpaceCast =
864 new AddrSpaceCastInst(I, NewPtrType, "", *InsertPoint);
865 AddrSpaceCast->setDebugLoc(I->getDebugLoc());
866 return AddrSpaceCast;
867 }
868 }
869
870 IRBuilder<> B(I);
871 if (NewPtrBits.getBitWidth() < MaskTy->getScalarSizeInBits()) {
872 MaskTy = MaskTy->getWithNewBitWidth(NewBitWidth: NewPtrBits.getBitWidth());
873 MaskOp = B.CreateTrunc(V: MaskOp, DestTy: MaskTy);
874 }
875 Value *NewPtr = operandWithNewAddressSpaceOrCreatePoison(
876 OperandUse: PtrOpUse, NewAddrSpace, ValueWithNewAddrSpace, PredicatedAS,
877 PoisonUsesToFix);
878 return B.CreateIntrinsic(ID: Intrinsic::ptrmask, OverloadTypes: {NewPtr->getType(), MaskTy},
879 Args: {NewPtr, MaskOp});
880}
881
882// Returns a clone of `I` with its operands converted to those specified in
883// ValueWithNewAddrSpace. Due to potential cycles in the data flow graph, an
884// operand whose address space needs to be modified might not exist in
885// ValueWithNewAddrSpace. In that case, uses poison as a placeholder operand and
886// adds that operand use to PoisonUsesToFix so that caller can fix them later.
887//
888// Note that we do not necessarily clone `I`, e.g., if it is an addrspacecast
889// from a pointer whose type already matches. Therefore, this function returns a
890// Value* instead of an Instruction*.
891Value *InferAddressSpacesImpl::cloneInstructionWithNewAddressSpace(
892 Instruction *I, unsigned NewAddrSpace,
893 const ValueToValueMapTy &ValueWithNewAddrSpace,
894 const PredicatedAddrSpaceMapTy &PredicatedAS,
895 SmallVectorImpl<const Use *> *PoisonUsesToFix) const {
896 Type *NewPtrType = getPtrOrVecOfPtrsWithNewAS(Ty: I->getType(), NewAddrSpace);
897
898 if (I->getOpcode() == Instruction::AddrSpaceCast) {
899 Value *Src = I->getOperand(i: 0);
900 // Because `I` is flat, the source address space must be specific.
901 // Therefore, the inferred address space must be the source space, according
902 // to our algorithm.
903 assert(Src->getType()->getPointerAddressSpace() == NewAddrSpace);
904 return Src;
905 }
906
907 if (IntrinsicInst *II = dyn_cast<IntrinsicInst>(Val: I)) {
908 // Technically the intrinsic ID is a pointer typed argument, so specially
909 // handle calls early.
910 assert(II->getIntrinsicID() == Intrinsic::ptrmask);
911 return clonePtrMaskWithNewAddressSpace(
912 I: II, NewAddrSpace, ValueWithNewAddrSpace, PredicatedAS, PoisonUsesToFix);
913 }
914
915 unsigned AS = TTI->getAssumedAddrSpace(V: I);
916 if (AS != UninitializedAddressSpace) {
917 // For the assumed address space, insert an `addrspacecast` to make that
918 // explicit.
919 Type *NewPtrTy = getPtrOrVecOfPtrsWithNewAS(Ty: I->getType(), NewAddrSpace: AS);
920 auto *NewI = new AddrSpaceCastInst(I, NewPtrTy);
921 NewI->insertAfter(InsertPos: I->getIterator());
922 NewI->setDebugLoc(I->getDebugLoc());
923 return NewI;
924 }
925
926 // Computes the converted pointer operands.
927 SmallVector<Value *, 4> NewPointerOperands;
928 for (const Use &OperandUse : I->operands()) {
929 if (!OperandUse.get()->getType()->isPtrOrPtrVectorTy())
930 NewPointerOperands.push_back(Elt: nullptr);
931 else
932 NewPointerOperands.push_back(Elt: operandWithNewAddressSpaceOrCreatePoison(
933 OperandUse, NewAddrSpace, ValueWithNewAddrSpace, PredicatedAS,
934 PoisonUsesToFix));
935 }
936
937 switch (I->getOpcode()) {
938 case Instruction::BitCast:
939 return new BitCastInst(NewPointerOperands[0], NewPtrType);
940 case Instruction::PHI: {
941 assert(I->getType()->isPtrOrPtrVectorTy());
942 PHINode *PHI = cast<PHINode>(Val: I);
943 PHINode *NewPHI = PHINode::Create(Ty: NewPtrType, NumReservedValues: PHI->getNumIncomingValues());
944 for (unsigned Index = 0; Index < PHI->getNumIncomingValues(); ++Index) {
945 unsigned OperandNo = PHINode::getOperandNumForIncomingValue(i: Index);
946 NewPHI->addIncoming(V: NewPointerOperands[OperandNo],
947 BB: PHI->getIncomingBlock(i: Index));
948 }
949 return NewPHI;
950 }
951 case Instruction::GetElementPtr: {
952 GetElementPtrInst *GEP = cast<GetElementPtrInst>(Val: I);
953 GetElementPtrInst *NewGEP = GetElementPtrInst::Create(
954 PointeeType: GEP->getSourceElementType(), Ptr: NewPointerOperands[0],
955 IdxList: SmallVector<Value *, 4>(GEP->indices()));
956 NewGEP->setIsInBounds(GEP->isInBounds());
957 return NewGEP;
958 }
959 case Instruction::Select:
960 assert(I->getType()->isPtrOrPtrVectorTy());
961 return SelectInst::Create(C: I->getOperand(i: 0), S1: NewPointerOperands[1],
962 S2: NewPointerOperands[2], NameStr: "", InsertBefore: nullptr, MDFrom: I);
963 case Instruction::IntToPtr: {
964 if (isNoopPtrIntCastPair(I2P: cast<Operator>(Val: I), DL: *DL, TTI)) {
965 Value *Src = cast<Operator>(Val: I->getOperand(i: 0))->getOperand(i: 0);
966 if (Src->getType() == NewPtrType)
967 return Src;
968
969 // If we had a no-op inttoptr/ptrtoint pair, we may still have inferred a
970 // source address space from a generic pointer source need to insert a
971 // cast back.
972 return new AddrSpaceCastInst(Src, NewPtrType);
973 }
974 assert(isSafeToCastIntToPtrAddrSpace(cast<Operator>(I)));
975 AddrSpaceCastInst *AsCast = new AddrSpaceCastInst(I, NewPtrType);
976 AsCast->insertAfter(InsertPos: I);
977 return AsCast;
978 }
979 default:
980 llvm_unreachable("Unexpected opcode");
981 }
982}
983
984// Similar to cloneInstructionWithNewAddressSpace, returns a clone of the
985// constant expression `CE` with its operands replaced as specified in
986// ValueWithNewAddrSpace.
987Value *InferAddressSpacesImpl::cloneConstantExprWithNewAddressSpace(
988 ConstantExpr *CE, unsigned NewAddrSpace,
989 const ValueToValueMapTy &ValueWithNewAddrSpace, const DataLayout *DL,
990 const TargetTransformInfo *TTI) const {
991 Type *TargetType =
992 CE->getType()->isPtrOrPtrVectorTy()
993 ? getPtrOrVecOfPtrsWithNewAS(Ty: CE->getType(), NewAddrSpace)
994 : CE->getType();
995
996 if (CE->getOpcode() == Instruction::AddrSpaceCast) {
997 // Because CE is flat, the source address space must be specific.
998 // Therefore, the inferred address space must be the source space according
999 // to our algorithm.
1000 assert(CE->getOperand(0)->getType()->getPointerAddressSpace() ==
1001 NewAddrSpace);
1002 return CE->getOperand(i_nocapture: 0);
1003 }
1004
1005 if (CE->getOpcode() == Instruction::BitCast) {
1006 if (Value *NewOperand = ValueWithNewAddrSpace.lookup(Val: CE->getOperand(i_nocapture: 0)))
1007 return ConstantExpr::getBitCast(C: cast<Constant>(Val: NewOperand), Ty: TargetType);
1008 return ConstantExpr::getAddrSpaceCast(C: CE, Ty: TargetType);
1009 }
1010
1011 if (CE->getOpcode() == Instruction::IntToPtr) {
1012 if (isNoopPtrIntCastPair(I2P: cast<Operator>(Val: CE), DL: *DL, TTI)) {
1013 Constant *Src = cast<ConstantExpr>(Val: CE->getOperand(i_nocapture: 0))->getOperand(i_nocapture: 0);
1014 assert(Src->getType()->getPointerAddressSpace() == NewAddrSpace);
1015 return Src;
1016 }
1017 assert(isSafeToCastIntToPtrAddrSpace(cast<Operator>(CE)));
1018 return ConstantExpr::getAddrSpaceCast(C: CE, Ty: TargetType);
1019 }
1020
1021 // Computes the operands of the new constant expression.
1022 bool IsNew = false;
1023 SmallVector<Constant *, 4> NewOperands;
1024 for (unsigned Index = 0; Index < CE->getNumOperands(); ++Index) {
1025 Constant *Operand = CE->getOperand(i_nocapture: Index);
1026 // If the address space of `Operand` needs to be modified, the new operand
1027 // with the new address space should already be in ValueWithNewAddrSpace
1028 // because (1) the constant expressions we consider (i.e. addrspacecast,
1029 // bitcast, and getelementptr) do not incur cycles in the data flow graph
1030 // and (2) this function is called on constant expressions in postorder.
1031 if (Value *NewOperand = ValueWithNewAddrSpace.lookup(Val: Operand)) {
1032 IsNew = true;
1033 NewOperands.push_back(Elt: cast<Constant>(Val: NewOperand));
1034 continue;
1035 }
1036 if (auto *CExpr = dyn_cast<ConstantExpr>(Val: Operand))
1037 if (Value *NewOperand = cloneConstantExprWithNewAddressSpace(
1038 CE: CExpr, NewAddrSpace, ValueWithNewAddrSpace, DL, TTI)) {
1039 IsNew = true;
1040 NewOperands.push_back(Elt: cast<Constant>(Val: NewOperand));
1041 continue;
1042 }
1043 // Otherwise, reuses the old operand.
1044 NewOperands.push_back(Elt: Operand);
1045 }
1046
1047 // If !IsNew, we will replace the Value with itself. However, replaced values
1048 // are assumed to wrapped in an addrspacecast cast later so drop it now.
1049 if (!IsNew)
1050 return nullptr;
1051
1052 if (CE->getOpcode() == Instruction::GetElementPtr) {
1053 // Needs to specify the source type while constructing a getelementptr
1054 // constant expression.
1055 return CE->getWithOperands(Ops: NewOperands, Ty: TargetType, /*OnlyIfReduced=*/false,
1056 SrcTy: cast<GEPOperator>(Val: CE)->getSourceElementType());
1057 }
1058
1059 return CE->getWithOperands(Ops: NewOperands, Ty: TargetType);
1060}
1061
1062// Returns a clone of the value `V`, with its operands replaced as specified in
1063// ValueWithNewAddrSpace. This function is called on every flat address
1064// expression whose address space needs to be modified, in postorder.
1065//
1066// See cloneInstructionWithNewAddressSpace for the meaning of PoisonUsesToFix.
1067Value *InferAddressSpacesImpl::cloneValueWithNewAddressSpace(
1068 Value *V, unsigned NewAddrSpace,
1069 const ValueToValueMapTy &ValueWithNewAddrSpace,
1070 const PredicatedAddrSpaceMapTy &PredicatedAS,
1071 SmallVectorImpl<const Use *> *PoisonUsesToFix) const {
1072 // All values in Postorder are flat address expressions.
1073 assert(V->getType()->getPointerAddressSpace() == FlatAddrSpace &&
1074 isAddressExpression(*V, *DL, TTI));
1075
1076 if (auto *Arg = dyn_cast<Argument>(Val: V)) {
1077 // Arguments are address space casted in the function body, as we do not
1078 // want to change the function signature.
1079 Function *F = Arg->getParent();
1080 BasicBlock::iterator Insert = F->getEntryBlock().getFirstNonPHIIt();
1081
1082 Type *NewPtrTy = PointerType::get(C&: Arg->getContext(), AddressSpace: NewAddrSpace);
1083 auto *NewI = new AddrSpaceCastInst(Arg, NewPtrTy);
1084 NewI->insertBefore(InsertPos: Insert);
1085 return NewI;
1086 }
1087
1088 if (Instruction *I = dyn_cast<Instruction>(Val: V)) {
1089 Value *NewV = cloneInstructionWithNewAddressSpace(
1090 I, NewAddrSpace, ValueWithNewAddrSpace, PredicatedAS, PoisonUsesToFix);
1091 if (Instruction *NewI = dyn_cast_or_null<Instruction>(Val: NewV)) {
1092 if (NewI->getParent() == nullptr) {
1093 NewI->insertBefore(InsertPos: I->getIterator());
1094 NewI->takeName(V: I);
1095 NewI->setDebugLoc(I->getDebugLoc());
1096 }
1097 }
1098 return NewV;
1099 }
1100
1101 return cloneConstantExprWithNewAddressSpace(
1102 CE: cast<ConstantExpr>(Val: V), NewAddrSpace, ValueWithNewAddrSpace, DL, TTI);
1103}
1104
1105// Defines the join operation on the address space lattice (see the file header
1106// comments).
1107unsigned InferAddressSpacesImpl::joinAddressSpaces(unsigned AS1,
1108 unsigned AS2) const {
1109 if (AS1 == AS2)
1110 return AS1;
1111
1112 if (AS1 == FlatAddrSpace || AS2 == FlatAddrSpace)
1113 return FlatAddrSpace;
1114
1115 if (AS1 == UninitializedAddressSpace)
1116 return AS2;
1117 if (AS2 == UninitializedAddressSpace)
1118 return AS1;
1119
1120 return TTI->getAddressSpaceJoin(AS1, AS2);
1121}
1122
1123bool InferAddressSpacesImpl::run(Function &CurFn) {
1124 F = &CurFn;
1125 DL = &F->getDataLayout();
1126 PtrIntCastPairs.clear();
1127
1128 if (AssumeDefaultIsFlatAddressSpace)
1129 FlatAddrSpace = 0;
1130
1131 if (FlatAddrSpace == UninitializedAddressSpace) {
1132 FlatAddrSpace = TTI->getFlatAddressSpace();
1133 if (FlatAddrSpace == UninitializedAddressSpace)
1134 return false;
1135 }
1136
1137 collectIntToPtrPointerOperand();
1138 // Collects all flat address expressions in postorder.
1139 std::vector<WeakTrackingVH> Postorder = collectFlatAddressExpressions(F&: *F);
1140
1141 // Runs a data-flow analysis to refine the address spaces of every expression
1142 // in Postorder.
1143 ValueToAddrSpaceMapTy InferredAddrSpace;
1144 PredicatedAddrSpaceMapTy PredicatedAS;
1145 inferAddressSpaces(Postorder, InferredAddrSpace, PredicatedAS);
1146
1147 // Changes the address spaces of the flat address expressions who are inferred
1148 // to point to a specific address space.
1149 return rewriteWithNewAddressSpaces(Postorder, InferredAddrSpace,
1150 PredicatedAS);
1151}
1152
1153void InferAddressSpacesImpl::enqueueUsers(
1154 Value &V, const ValueToAddrSpaceMapTy &InferredAddrSpace,
1155 SetVector<Value *> &Worklist) const {
1156 for (Value *User : V.users()) {
1157 // Skip if User is already in the worklist.
1158 if (Worklist.count(key: User))
1159 continue;
1160
1161 ValueToAddrSpaceMapTy::const_iterator Pos = InferredAddrSpace.find(Val: User);
1162 // Our algorithm only updates the address spaces of flat address
1163 // expressions, which are those in InferredAddrSpace.
1164 if (Pos == InferredAddrSpace.end())
1165 continue;
1166
1167 // Function updateAddressSpace moves the address space down a lattice path.
1168 // Therefore, nothing to do if User is already inferred as flat (the bottom
1169 // element in the lattice).
1170 if (Pos->second == FlatAddrSpace)
1171 continue;
1172
1173 Worklist.insert(X: User);
1174 }
1175}
1176
1177void InferAddressSpacesImpl::runToFixPoint(
1178 SetVector<Value *> &Worklist, ValueToAddrSpaceMapTy &InferredAddrSpace,
1179 PredicatedAddrSpaceMapTy &PredicatedAS) const {
1180 while (!Worklist.empty()) {
1181 Value *V = Worklist.pop_back_val();
1182
1183 // Try to update the address space of the stack top according to the
1184 // address spaces of its operands.
1185 if (!updateAddressSpace(V: *V, InferredAddrSpace, PredicatedAS))
1186 continue;
1187
1188 enqueueUsers(V&: *V, InferredAddrSpace, Worklist);
1189 }
1190}
1191
1192// Constants need to be tracked through RAUW to handle cases with nested
1193// constant expressions, so wrap values in WeakTrackingVH.
1194void InferAddressSpacesImpl::inferAddressSpaces(
1195 ArrayRef<WeakTrackingVH> Postorder,
1196 ValueToAddrSpaceMapTy &InferredAddrSpace,
1197 PredicatedAddrSpaceMapTy &PredicatedAS) const {
1198 SetVector<Value *> Worklist(llvm::from_range, Postorder);
1199 // Initially, all expressions are in the uninitialized address space.
1200 for (Value *V : Postorder)
1201 InferredAddrSpace[V] = UninitializedAddressSpace;
1202
1203 runToFixPoint(Worklist, InferredAddrSpace, PredicatedAS);
1204
1205 // A value still uninitialized here is stuck in a cycle of uninitialized
1206 // values and carries no address space information. Lower it to flat so its
1207 // users join to flat, instead of being rewritten to reference an operand
1208 // that rewriteWithNewAddressSpaces() never converts.
1209 SmallVector<Value *, 4> Lowered;
1210 for (Value *V : Postorder) {
1211 ValueToAddrSpaceMapTy::iterator I = InferredAddrSpace.find(Val: V);
1212 if (I->second == UninitializedAddressSpace) {
1213 I->second = FlatAddrSpace;
1214 Lowered.push_back(Elt: V);
1215 }
1216 }
1217
1218 for (Value *V : Lowered)
1219 enqueueUsers(V&: *V, InferredAddrSpace, Worklist);
1220
1221 runToFixPoint(Worklist, InferredAddrSpace, PredicatedAS);
1222}
1223
1224unsigned
1225InferAddressSpacesImpl::getPredicatedAddrSpace(const Value &Ptr,
1226 const Value *UserCtx) const {
1227 const Instruction *UserCtxI = dyn_cast<Instruction>(Val: UserCtx);
1228 if (!UserCtxI)
1229 return UninitializedAddressSpace;
1230
1231 const Value *StrippedPtr = Ptr.stripInBoundsOffsets();
1232 for (auto &AssumeVH : AC.assumptionsFor(V: StrippedPtr)) {
1233 if (!AssumeVH)
1234 continue;
1235 CallInst *CI = cast<CallInst>(Val&: AssumeVH);
1236 if (!isValidAssumeForContext(I: CI, CtxI: UserCtxI, DT))
1237 continue;
1238
1239 const Value *Ptr;
1240 unsigned AS;
1241 std::tie(args&: Ptr, args&: AS) = TTI->getPredicatedAddrSpace(V: CI->getArgOperand(i: 0));
1242 if (Ptr)
1243 return AS;
1244 }
1245
1246 return UninitializedAddressSpace;
1247}
1248
1249bool InferAddressSpacesImpl::updateAddressSpace(
1250 const Value &V, ValueToAddrSpaceMapTy &InferredAddrSpace,
1251 PredicatedAddrSpaceMapTy &PredicatedAS) const {
1252 assert(InferredAddrSpace.count(&V));
1253
1254 LLVM_DEBUG(dbgs() << "Updating the address space of\n " << V << '\n');
1255
1256 // The new inferred address space equals the join of the address spaces
1257 // of all its pointer operands.
1258 unsigned NewAS = UninitializedAddressSpace;
1259
1260 // isAddressExpression should guarantee that V is an operator or an argument.
1261 assert(isa<Operator>(V) || isa<Argument>(V));
1262
1263 unsigned AS = TTI->getAssumedAddrSpace(V: &V);
1264 if (AS != UninitializedAddressSpace) {
1265 // Use the assumed address space directly.
1266 NewAS = AS;
1267 } else {
1268 // Otherwise, infer the address space from its pointer operands.
1269 SmallVector<Constant *, 2> ConstantPtrOps;
1270 SmallVector<Value *, 2> PtrOps = getPointerOperands(V, DL: *DL, TTI);
1271 for (Value *PtrOperand : PtrOps) {
1272 auto I = InferredAddrSpace.find(Val: PtrOperand);
1273 unsigned OperandAS;
1274 if (I == InferredAddrSpace.end()) {
1275 OperandAS = PtrOperand->getType()->getPointerAddressSpace();
1276 if (auto *C = dyn_cast<Constant>(Val: PtrOperand);
1277 C && OperandAS == FlatAddrSpace) {
1278 // Defer joining the address space of constant pointer operands.
1279 ConstantPtrOps.push_back(Elt: C);
1280 continue;
1281 }
1282 if (OperandAS == FlatAddrSpace) {
1283 // Check AC for assumption dominating V.
1284 unsigned AS = getPredicatedAddrSpace(Ptr: *PtrOperand, UserCtx: &V);
1285 if (AS != UninitializedAddressSpace) {
1286 LLVM_DEBUG(dbgs()
1287 << " deduce operand AS from the predicate addrspace "
1288 << AS << '\n');
1289 OperandAS = AS;
1290 // Record this use with the predicated AS.
1291 PredicatedAS[std::make_pair(x: &V, y&: PtrOperand)] = OperandAS;
1292 }
1293 }
1294 } else
1295 OperandAS = I->second;
1296
1297 // join(flat, *) = flat. So we can break if NewAS is already flat.
1298 NewAS = joinAddressSpaces(AS1: NewAS, AS2: OperandAS);
1299 if (NewAS == FlatAddrSpace)
1300 break;
1301 }
1302
1303 if (NewAS != FlatAddrSpace && NewAS != UninitializedAddressSpace) {
1304 if (any_of(Range&: ConstantPtrOps, P: [=](Constant *C) {
1305 return !isSafeToCastConstAddrSpace(C, NewAS);
1306 }))
1307 NewAS = FlatAddrSpace;
1308 }
1309
1310 // operator(flat const, flat const, ...) -> flat
1311 if (NewAS == UninitializedAddressSpace &&
1312 PtrOps.size() == ConstantPtrOps.size())
1313 NewAS = FlatAddrSpace;
1314 }
1315
1316 unsigned OldAS = InferredAddrSpace.lookup(Val: &V);
1317 assert(OldAS != FlatAddrSpace);
1318 if (OldAS == NewAS)
1319 return false;
1320
1321 // If any updates are made, grabs its users to the worklist because
1322 // their address spaces can also be possibly updated.
1323 LLVM_DEBUG(dbgs() << " to " << NewAS << '\n');
1324 InferredAddrSpace[&V] = NewAS;
1325 return true;
1326}
1327
1328/// Replace operand \p OpIdx in \p Inst, if the value is the same as \p OldVal
1329/// with \p NewVal.
1330static bool replaceOperandIfSame(Instruction *Inst, unsigned OpIdx,
1331 Value *OldVal, Value *NewVal) {
1332 Use &U = Inst->getOperandUse(i: OpIdx);
1333 if (U.get() == OldVal) {
1334 U.set(NewVal);
1335 return true;
1336 }
1337
1338 return false;
1339}
1340
1341template <typename InstrType>
1342static bool replaceSimplePointerUse(const TargetTransformInfo &TTI,
1343 InstrType *MemInstr, Value *OldV,
1344 Value *NewV) {
1345 unsigned AddrSpace = NewV->getType()->getPointerAddressSpace();
1346 if (!MemInstr->isVolatile() || TTI.hasVolatileVariant(I: MemInstr, AddrSpace)) {
1347 return replaceOperandIfSame(MemInstr, InstrType::getPointerOperandIndex(),
1348 OldV, NewV);
1349 }
1350
1351 return false;
1352}
1353
1354/// If \p OldV is used as the pointer operand of a compatible memory operation
1355/// \p Inst, replaces the pointer operand with NewV.
1356///
1357/// This covers memory instructions with a single pointer operand that can have
1358/// its address space changed by simply mutating the use to a new value.
1359///
1360/// \p returns true the user replacement was made.
1361static bool replaceIfSimplePointerUse(const TargetTransformInfo &TTI,
1362 User *Inst, Value *OldV, Value *NewV) {
1363 if (auto *LI = dyn_cast<LoadInst>(Val: Inst))
1364 return replaceSimplePointerUse(TTI, MemInstr: LI, OldV, NewV);
1365
1366 if (auto *SI = dyn_cast<StoreInst>(Val: Inst))
1367 return replaceSimplePointerUse(TTI, MemInstr: SI, OldV, NewV);
1368
1369 if (auto *RMW = dyn_cast<AtomicRMWInst>(Val: Inst))
1370 return replaceSimplePointerUse(TTI, MemInstr: RMW, OldV, NewV);
1371
1372 if (auto *CmpX = dyn_cast<AtomicCmpXchgInst>(Val: Inst))
1373 return replaceSimplePointerUse(TTI, MemInstr: CmpX, OldV, NewV);
1374
1375 return false;
1376}
1377
1378/// Update memory intrinsic uses that require more complex processing than
1379/// simple memory instructions. These require re-mangling and may have multiple
1380/// pointer operands.
1381static bool handleMemIntrinsicPtrUse(MemIntrinsic *MI, Value *OldV,
1382 Value *NewV) {
1383 IRBuilder<> B(MI);
1384 if (auto *MSI = dyn_cast<MemSetInst>(Val: MI)) {
1385 B.CreateMemSet(Ptr: NewV, Val: MSI->getValue(), Size: MSI->getLength(), Align: MSI->getDestAlign(),
1386 isVolatile: false, // isVolatile
1387 AAInfo: MI->getAAMetadata());
1388 } else if (auto *MTI = dyn_cast<MemTransferInst>(Val: MI)) {
1389 Value *Src = MTI->getRawSource();
1390 Value *Dest = MTI->getRawDest();
1391
1392 // Be careful in case this is a self-to-self copy.
1393 if (Src == OldV)
1394 Src = NewV;
1395
1396 if (Dest == OldV)
1397 Dest = NewV;
1398
1399 if (auto *MCI = dyn_cast<MemCpyInst>(Val: MTI)) {
1400 if (MCI->isForceInlined())
1401 B.CreateMemCpyInline(Dst: Dest, DstAlign: MTI->getDestAlign(), Src,
1402 SrcAlign: MTI->getSourceAlign(), Size: MTI->getLength(),
1403 isVolatile: false, // isVolatile
1404 AAInfo: MI->getAAMetadata());
1405 else
1406 B.CreateMemCpy(Dst: Dest, DstAlign: MTI->getDestAlign(), Src, SrcAlign: MTI->getSourceAlign(),
1407 Size: MTI->getLength(),
1408 isVolatile: false, // isVolatile
1409 AAInfo: MI->getAAMetadata());
1410 } else {
1411 assert(isa<MemMoveInst>(MTI));
1412 B.CreateMemMove(Dst: Dest, DstAlign: MTI->getDestAlign(), Src, SrcAlign: MTI->getSourceAlign(),
1413 Size: MTI->getLength(),
1414 isVolatile: false, // isVolatile
1415 AAInfo: MI->getAAMetadata());
1416 }
1417 } else
1418 llvm_unreachable("unhandled MemIntrinsic");
1419
1420 MI->eraseFromParent();
1421 return true;
1422}
1423
1424// \p returns true if it is OK to change the address space of constant \p C with
1425// a ConstantExpr addrspacecast.
1426bool InferAddressSpacesImpl::isSafeToCastConstAddrSpace(Constant *C,
1427 unsigned NewAS) const {
1428 assert(NewAS != UninitializedAddressSpace);
1429
1430 unsigned SrcAS = C->getType()->getPointerAddressSpace();
1431 if (SrcAS == NewAS || isa<UndefValue>(Val: C))
1432 return true;
1433
1434 // Prevent illegal casts between different non-flat address spaces.
1435 if (SrcAS != FlatAddrSpace && NewAS != FlatAddrSpace)
1436 return false;
1437
1438 if (isa<ConstantPointerNull>(Val: C) || isa<ConstantAggregateZero>(Val: C))
1439 return true;
1440
1441 if (auto *Op = dyn_cast<Operator>(Val: C)) {
1442 // If we already have a constant addrspacecast, it should be safe to cast it
1443 // off.
1444 if (Op->getOpcode() == Instruction::AddrSpaceCast)
1445 return isSafeToCastConstAddrSpace(C: cast<Constant>(Val: Op->getOperand(i: 0)),
1446 NewAS);
1447
1448 if (Op->getOpcode() == Instruction::IntToPtr &&
1449 Op->getType()->getPointerAddressSpace() == FlatAddrSpace)
1450 return true;
1451 }
1452
1453 return false;
1454}
1455
1456static Value::use_iterator skipToNextUser(Value::use_iterator I,
1457 Value::use_iterator End) {
1458 User *CurUser = I->getUser();
1459 ++I;
1460
1461 while (I != End && I->getUser() == CurUser)
1462 ++I;
1463
1464 return I;
1465}
1466
1467void InferAddressSpacesImpl::performPointerReplacement(
1468 Value *V, Value *NewV, Use &U, ValueToValueMapTy &ValueWithNewAddrSpace,
1469 SmallVectorImpl<Instruction *> &DeadInstructions) const {
1470
1471 User *CurUser = U.getUser();
1472
1473 if (replaceIfSimplePointerUse(TTI: *TTI, Inst: CurUser, OldV: V, NewV))
1474 return;
1475
1476 // Skip if the current user is the new value itself.
1477 if (CurUser == NewV)
1478 return;
1479
1480 auto *CurUserI = dyn_cast<Instruction>(Val: CurUser);
1481 if (!CurUserI || CurUserI->getFunction() != F)
1482 return;
1483
1484 // Handle more complex cases like intrinsic that need to be remangled.
1485 if (auto *MI = dyn_cast<MemIntrinsic>(Val: CurUser)) {
1486 if (!MI->isVolatile() && handleMemIntrinsicPtrUse(MI, OldV: V, NewV))
1487 return;
1488 }
1489
1490 if (auto *II = dyn_cast<IntrinsicInst>(Val: CurUser)) {
1491 if (rewriteIntrinsicOperands(II, OldV: V, NewV))
1492 return;
1493 }
1494
1495 if (ICmpInst *Cmp = dyn_cast<ICmpInst>(Val: CurUserI)) {
1496 // If we can infer that both pointers are in the same addrspace,
1497 // transform e.g.
1498 // %cmp = icmp eq float* %p, %q
1499 // into
1500 // %cmp = icmp eq float addrspace(3)* %new_p, %new_q
1501
1502 unsigned NewAS = NewV->getType()->getPointerAddressSpace();
1503 int SrcIdx = U.getOperandNo();
1504 int OtherIdx = (SrcIdx == 0) ? 1 : 0;
1505 Value *OtherSrc = Cmp->getOperand(i_nocapture: OtherIdx);
1506
1507 if (Value *OtherNewV = ValueWithNewAddrSpace.lookup(Val: OtherSrc)) {
1508 if (OtherNewV->getType()->getPointerAddressSpace() == NewAS) {
1509 Cmp->setOperand(i_nocapture: OtherIdx, Val_nocapture: OtherNewV);
1510 Cmp->setOperand(i_nocapture: SrcIdx, Val_nocapture: NewV);
1511 return;
1512 }
1513 }
1514
1515 // Even if the type mismatches, we can cast the constant.
1516 if (auto *KOtherSrc = dyn_cast<Constant>(Val: OtherSrc)) {
1517 if (isSafeToCastConstAddrSpace(C: KOtherSrc, NewAS)) {
1518 Cmp->setOperand(i_nocapture: SrcIdx, Val_nocapture: NewV);
1519 Cmp->setOperand(i_nocapture: OtherIdx, Val_nocapture: ConstantExpr::getAddrSpaceCast(
1520 C: KOtherSrc, Ty: NewV->getType()));
1521 return;
1522 }
1523 }
1524 }
1525
1526 if (AddrSpaceCastInst *ASC = dyn_cast<AddrSpaceCastInst>(Val: CurUserI)) {
1527 unsigned NewAS = NewV->getType()->getPointerAddressSpace();
1528 if (ASC->getDestAddressSpace() == NewAS) {
1529 ASC->replaceAllUsesWith(V: NewV);
1530 DeadInstructions.push_back(Elt: ASC);
1531 return;
1532 }
1533 }
1534
1535 // Otherwise, replaces the use with flat(NewV).
1536 if (isa<Instruction>(Val: V) || isa<Instruction>(Val: NewV)) {
1537 // Don't create a copy of the original addrspacecast.
1538 if (U == V && isa<AddrSpaceCastInst>(Val: V))
1539 return;
1540
1541 // Insert the addrspacecast after NewV.
1542 BasicBlock::iterator InsertPos;
1543 if (Instruction *NewVInst = dyn_cast<Instruction>(Val: NewV))
1544 InsertPos = std::next(x: NewVInst->getIterator());
1545 else
1546 InsertPos = std::next(x: cast<Instruction>(Val: V)->getIterator());
1547
1548 while (isa<PHINode>(Val: InsertPos))
1549 ++InsertPos;
1550 // This instruction may contain multiple uses of V, update them all.
1551 CurUser->replaceUsesOfWith(
1552 From: V, To: new AddrSpaceCastInst(NewV, V->getType(), "", InsertPos));
1553 } else {
1554 CurUserI->replaceUsesOfWith(
1555 From: V, To: ConstantExpr::getAddrSpaceCast(C: cast<Constant>(Val: NewV), Ty: V->getType()));
1556 }
1557}
1558
1559bool InferAddressSpacesImpl::rewriteWithNewAddressSpaces(
1560 ArrayRef<WeakTrackingVH> Postorder,
1561 const ValueToAddrSpaceMapTy &InferredAddrSpace,
1562 const PredicatedAddrSpaceMapTy &PredicatedAS) const {
1563 // For each address expression to be modified, creates a clone of it with its
1564 // pointer operands converted to the new address space. Since the pointer
1565 // operands are converted, the clone is naturally in the new address space by
1566 // construction.
1567 ValueToValueMapTy ValueWithNewAddrSpace;
1568 SmallVector<const Use *, 32> PoisonUsesToFix;
1569 for (Value *V : Postorder) {
1570 unsigned NewAddrSpace = InferredAddrSpace.lookup(Val: V);
1571
1572 // In some degenerate cases (e.g. invalid IR in unreachable code), we may
1573 // not even infer the value to have its original address space.
1574 if (NewAddrSpace == UninitializedAddressSpace)
1575 continue;
1576
1577 if (V->getType()->getPointerAddressSpace() != NewAddrSpace) {
1578 Value *New =
1579 cloneValueWithNewAddressSpace(V, NewAddrSpace, ValueWithNewAddrSpace,
1580 PredicatedAS, PoisonUsesToFix: &PoisonUsesToFix);
1581 if (New)
1582 ValueWithNewAddrSpace[V] = New;
1583 }
1584 }
1585
1586 if (ValueWithNewAddrSpace.empty())
1587 return false;
1588
1589 // Fixes all the poison uses generated by cloneInstructionWithNewAddressSpace.
1590 for (const Use *PoisonUse : PoisonUsesToFix) {
1591 User *V = PoisonUse->getUser();
1592 User *NewV = cast_or_null<User>(Val: ValueWithNewAddrSpace.lookup(Val: V));
1593 if (!NewV)
1594 continue;
1595
1596 unsigned OperandNo = PoisonUse->getOperandNo();
1597 assert(isa<PoisonValue>(NewV->getOperand(OperandNo)));
1598 unsigned NewAS =
1599 NewV->getOperand(i: OperandNo)->getType()->getPointerAddressSpace();
1600 Value *NewOp = operandWithNewAddressSpaceOrCreatePoison(
1601 OperandUse: *PoisonUse, NewAddrSpace: NewAS, ValueWithNewAddrSpace, PredicatedAS, PoisonUsesToFix: nullptr);
1602 NewV->setOperand(i: OperandNo, Val: NewOp);
1603 }
1604
1605 SmallVector<Instruction *, 16> DeadInstructions;
1606 ValueToValueMapTy VMap;
1607 ValueMapper VMapper(VMap, RF_NoModuleLevelChanges | RF_IgnoreMissingLocals);
1608
1609 // Replaces the uses of the old address expressions with the new ones.
1610 for (const WeakTrackingVH &WVH : Postorder) {
1611 assert(WVH && "value was unexpectedly deleted");
1612 Value *V = WVH;
1613 Value *NewV = ValueWithNewAddrSpace.lookup(Val: V);
1614 if (NewV == nullptr)
1615 continue;
1616
1617 LLVM_DEBUG(dbgs() << "Replacing the uses of " << *V << "\n with\n "
1618 << *NewV << '\n');
1619
1620 if (Constant *C = dyn_cast<Constant>(Val: V)) {
1621 Constant *Replace =
1622 ConstantExpr::getAddrSpaceCast(C: cast<Constant>(Val: NewV), Ty: C->getType());
1623 if (C != Replace) {
1624 LLVM_DEBUG(dbgs() << "Inserting replacement const cast: " << Replace
1625 << ": " << *Replace << '\n');
1626 SmallVector<User *, 16> WorkList;
1627 for (User *U : make_early_inc_range(Range: C->users())) {
1628 if (auto *I = dyn_cast<Instruction>(Val: U)) {
1629 if (I->getFunction() == F)
1630 I->replaceUsesOfWith(From: C, To: Replace);
1631 } else {
1632 WorkList.append(in_start: U->user_begin(), in_end: U->user_end());
1633 }
1634 }
1635 if (!WorkList.empty()) {
1636 VMap[C] = Replace;
1637 DenseSet<User *> Visited{WorkList.begin(), WorkList.end()};
1638 while (!WorkList.empty()) {
1639 User *U = WorkList.pop_back_val();
1640 if (auto *I = dyn_cast<Instruction>(Val: U)) {
1641 if (I->getFunction() == F)
1642 VMapper.remapInstruction(I&: *I);
1643 continue;
1644 }
1645 for (User *U2 : U->users())
1646 if (Visited.insert(V: U2).second)
1647 WorkList.push_back(Elt: U2);
1648 }
1649 }
1650 V = Replace;
1651 }
1652 }
1653
1654 Value::use_iterator I, E, Next;
1655 for (I = V->use_begin(), E = V->use_end(); I != E;) {
1656 Use &U = *I;
1657
1658 // Some users may see the same pointer operand in multiple operands. Skip
1659 // to the next instruction.
1660 I = skipToNextUser(I, End: E);
1661
1662 performPointerReplacement(V, NewV, U, ValueWithNewAddrSpace,
1663 DeadInstructions);
1664 }
1665
1666 if (V->use_empty()) {
1667 if (Instruction *I = dyn_cast<Instruction>(Val: V))
1668 DeadInstructions.push_back(Elt: I);
1669 }
1670 }
1671
1672 // Deleting one instruction may recursively delete another queued
1673 // instruction. Create handles before the first deletion so overlapping
1674 // entries are nulled instead of leaving dangling pointers.
1675 auto DeadInstructionHandles =
1676 to_vector_of<WeakTrackingVH, 16>(Range&: DeadInstructions);
1677 RecursivelyDeleteTriviallyDeadInstructions(DeadInsts&: DeadInstructionHandles);
1678
1679 return true;
1680}
1681
1682bool InferAddressSpaces::runOnFunction(Function &F) {
1683 if (skipFunction(F))
1684 return false;
1685
1686 auto *DTWP = getAnalysisIfAvailable<DominatorTreeWrapperPass>();
1687 DominatorTree *DT = DTWP ? &DTWP->getDomTree() : nullptr;
1688 return InferAddressSpacesImpl(
1689 getAnalysis<AssumptionCacheTracker>().getAssumptionCache(F), DT,
1690 &getAnalysis<TargetTransformInfoWrapperPass>().getTTI(F),
1691 FlatAddrSpace, /*AssumeDefaultIsFlatAddressSpace=*/false)
1692 .run(CurFn&: F);
1693}
1694
1695FunctionPass *llvm::createInferAddressSpacesPass(unsigned AddressSpace) {
1696 return new InferAddressSpaces(AddressSpace);
1697}
1698
1699InferAddressSpacesPass::InferAddressSpacesPass(
1700 bool AssumeDefaultIsFlatAddressSpace)
1701 : FlatAddrSpace(UninitializedAddressSpace),
1702 AssumeDefaultIsFlatAddressSpace(AssumeDefaultIsFlatAddressSpace) {}
1703InferAddressSpacesPass::InferAddressSpacesPass(
1704 unsigned AddressSpace, bool AssumeDefaultIsFlatAddressSpace)
1705 : FlatAddrSpace(AddressSpace),
1706 AssumeDefaultIsFlatAddressSpace(AssumeDefaultIsFlatAddressSpace) {}
1707
1708PreservedAnalyses InferAddressSpacesPass::run(Function &F,
1709 FunctionAnalysisManager &AM) {
1710 bool Changed =
1711 InferAddressSpacesImpl(AM.getResult<AssumptionAnalysis>(IR&: F),
1712 AM.getCachedResult<DominatorTreeAnalysis>(IR&: F),
1713 &AM.getResult<TargetIRAnalysis>(IR&: F), FlatAddrSpace,
1714 AssumeDefaultIsFlatAddressSpace)
1715 .run(CurFn&: F);
1716 if (Changed) {
1717 PreservedAnalyses PA;
1718 PA.preserveSet<CFGAnalyses>();
1719 return PA;
1720 }
1721 return PreservedAnalyses::all();
1722}
1723
1724void InferAddressSpacesPass::printPipeline(
1725 raw_ostream &OS, function_ref<StringRef(StringRef)> MapClassName2PassName) {
1726 static_cast<PassInfoMixin<InferAddressSpacesPass> *>(this)->printPipeline(
1727 OS, MapClassName2PassName);
1728 if (AssumeDefaultIsFlatAddressSpace)
1729 OS << "<assume-default-is-flat-addrspace>";
1730}
1731