LLVM 18.0.0git
InferAddressSpaces.cpp
Go to the documentation of this file.
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 the generic 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
92#include "llvm/ADT/ArrayRef.h"
93#include "llvm/ADT/DenseMap.h"
94#include "llvm/ADT/DenseSet.h"
95#include "llvm/ADT/SetVector.h"
100#include "llvm/IR/BasicBlock.h"
101#include "llvm/IR/Constant.h"
102#include "llvm/IR/Constants.h"
103#include "llvm/IR/Dominators.h"
104#include "llvm/IR/Function.h"
105#include "llvm/IR/IRBuilder.h"
106#include "llvm/IR/InstIterator.h"
107#include "llvm/IR/Instruction.h"
108#include "llvm/IR/Instructions.h"
110#include "llvm/IR/Intrinsics.h"
111#include "llvm/IR/LLVMContext.h"
112#include "llvm/IR/Operator.h"
113#include "llvm/IR/PassManager.h"
114#include "llvm/IR/Type.h"
115#include "llvm/IR/Use.h"
116#include "llvm/IR/User.h"
117#include "llvm/IR/Value.h"
118#include "llvm/IR/ValueHandle.h"
120#include "llvm/Pass.h"
121#include "llvm/Support/Casting.h"
124#include "llvm/Support/Debug.h"
130#include <cassert>
131#include <iterator>
132#include <limits>
133#include <utility>
134#include <vector>
135
136#define DEBUG_TYPE "infer-address-spaces"
137
138using namespace llvm;
139
141 "assume-default-is-flat-addrspace", cl::init(false), cl::ReallyHidden,
142 cl::desc("The default address space is assumed as the flat address space. "
143 "This is mainly for test purpose."));
144
145static const unsigned UninitializedAddressSpace =
146 std::numeric_limits<unsigned>::max();
147
148namespace {
149
150using ValueToAddrSpaceMapTy = DenseMap<const Value *, unsigned>;
151// Different from ValueToAddrSpaceMapTy, where a new addrspace is inferred on
152// the *def* of a value, PredicatedAddrSpaceMapTy is map where a new
153// addrspace is inferred on the *use* of a pointer. This map is introduced to
154// infer addrspace from the addrspace predicate assumption built from assume
155// intrinsic. In that scenario, only specific uses (under valid assumption
156// context) could be inferred with a new addrspace.
157using PredicatedAddrSpaceMapTy =
159using PostorderStackTy = llvm::SmallVector<PointerIntPair<Value *, 1, bool>, 4>;
160
161class InferAddressSpaces : public FunctionPass {
162 unsigned FlatAddrSpace = 0;
163
164public:
165 static char ID;
166
167 InferAddressSpaces()
168 : FunctionPass(ID), FlatAddrSpace(UninitializedAddressSpace) {
170 }
171 InferAddressSpaces(unsigned AS) : FunctionPass(ID), FlatAddrSpace(AS) {
173 }
174
175 void getAnalysisUsage(AnalysisUsage &AU) const override {
176 AU.setPreservesCFG();
180 }
181
182 bool runOnFunction(Function &F) override;
183};
184
185class InferAddressSpacesImpl {
186 AssumptionCache &AC;
187 const DominatorTree *DT = nullptr;
188 const TargetTransformInfo *TTI = nullptr;
189 const DataLayout *DL = nullptr;
190
191 /// Target specific address space which uses of should be replaced if
192 /// possible.
193 unsigned FlatAddrSpace = 0;
194
195 // Try to update the address space of V. If V is updated, returns true and
196 // false otherwise.
197 bool updateAddressSpace(const Value &V,
198 ValueToAddrSpaceMapTy &InferredAddrSpace,
199 PredicatedAddrSpaceMapTy &PredicatedAS) const;
200
201 // Tries to infer the specific address space of each address expression in
202 // Postorder.
203 void inferAddressSpaces(ArrayRef<WeakTrackingVH> Postorder,
204 ValueToAddrSpaceMapTy &InferredAddrSpace,
205 PredicatedAddrSpaceMapTy &PredicatedAS) const;
206
207 bool isSafeToCastConstAddrSpace(Constant *C, unsigned NewAS) const;
208
209 Value *cloneInstructionWithNewAddressSpace(
210 Instruction *I, unsigned NewAddrSpace,
211 const ValueToValueMapTy &ValueWithNewAddrSpace,
212 const PredicatedAddrSpaceMapTy &PredicatedAS,
213 SmallVectorImpl<const Use *> *PoisonUsesToFix) const;
214
215 // Changes the flat address expressions in function F to point to specific
216 // address spaces if InferredAddrSpace says so. Postorder is the postorder of
217 // all flat expressions in the use-def graph of function F.
218 bool
219 rewriteWithNewAddressSpaces(ArrayRef<WeakTrackingVH> Postorder,
220 const ValueToAddrSpaceMapTy &InferredAddrSpace,
221 const PredicatedAddrSpaceMapTy &PredicatedAS,
222 Function *F) const;
223
224 void appendsFlatAddressExpressionToPostorderStack(
225 Value *V, PostorderStackTy &PostorderStack,
226 DenseSet<Value *> &Visited) const;
227
228 bool rewriteIntrinsicOperands(IntrinsicInst *II, Value *OldV,
229 Value *NewV) const;
230 void collectRewritableIntrinsicOperands(IntrinsicInst *II,
231 PostorderStackTy &PostorderStack,
232 DenseSet<Value *> &Visited) const;
233
234 std::vector<WeakTrackingVH> collectFlatAddressExpressions(Function &F) const;
235
236 Value *cloneValueWithNewAddressSpace(
237 Value *V, unsigned NewAddrSpace,
238 const ValueToValueMapTy &ValueWithNewAddrSpace,
239 const PredicatedAddrSpaceMapTy &PredicatedAS,
240 SmallVectorImpl<const Use *> *PoisonUsesToFix) const;
241 unsigned joinAddressSpaces(unsigned AS1, unsigned AS2) const;
242
243 unsigned getPredicatedAddrSpace(const Value &V, Value *Opnd) const;
244
245public:
246 InferAddressSpacesImpl(AssumptionCache &AC, const DominatorTree *DT,
247 const TargetTransformInfo *TTI, unsigned FlatAddrSpace)
248 : AC(AC), DT(DT), TTI(TTI), FlatAddrSpace(FlatAddrSpace) {}
249 bool run(Function &F);
250};
251
252} // end anonymous namespace
253
254char InferAddressSpaces::ID = 0;
255
256INITIALIZE_PASS_BEGIN(InferAddressSpaces, DEBUG_TYPE, "Infer address spaces",
257 false, false)
260INITIALIZE_PASS_END(InferAddressSpaces, DEBUG_TYPE, "Infer address spaces",
262
263static Type *getPtrOrVecOfPtrsWithNewAS(Type *Ty, unsigned NewAddrSpace) {
264 assert(Ty->isPtrOrPtrVectorTy());
265 PointerType *NPT = PointerType::get(Ty->getContext(), NewAddrSpace);
266 return Ty->getWithNewType(NPT);
267}
268
269// Check whether that's no-op pointer bicast using a pair of
270// `ptrtoint`/`inttoptr` due to the missing no-op pointer bitcast over
271// different address spaces.
272static bool isNoopPtrIntCastPair(const Operator *I2P, const DataLayout &DL,
273 const TargetTransformInfo *TTI) {
274 assert(I2P->getOpcode() == Instruction::IntToPtr);
275 auto *P2I = dyn_cast<Operator>(I2P->getOperand(0));
276 if (!P2I || P2I->getOpcode() != Instruction::PtrToInt)
277 return false;
278 // Check it's really safe to treat that pair of `ptrtoint`/`inttoptr` as a
279 // no-op cast. Besides checking both of them are no-op casts, as the
280 // reinterpreted pointer may be used in other pointer arithmetic, we also
281 // need to double-check that through the target-specific hook. That ensures
282 // the underlying target also agrees that's a no-op address space cast and
283 // pointer bits are preserved.
284 // The current IR spec doesn't have clear rules on address space casts,
285 // especially a clear definition for pointer bits in non-default address
286 // spaces. It would be undefined if that pointer is dereferenced after an
287 // invalid reinterpret cast. Also, due to the unclearness for the meaning of
288 // bits in non-default address spaces in the current spec, the pointer
289 // arithmetic may also be undefined after invalid pointer reinterpret cast.
290 // However, as we confirm through the target hooks that it's a no-op
291 // addrspacecast, it doesn't matter since the bits should be the same.
292 unsigned P2IOp0AS = P2I->getOperand(0)->getType()->getPointerAddressSpace();
293 unsigned I2PAS = I2P->getType()->getPointerAddressSpace();
295 I2P->getOperand(0)->getType(), I2P->getType(),
296 DL) &&
298 P2I->getOperand(0)->getType(), P2I->getType(),
299 DL) &&
300 (P2IOp0AS == I2PAS || TTI->isNoopAddrSpaceCast(P2IOp0AS, I2PAS));
301}
302
303// Returns true if V is an address expression.
304// TODO: Currently, we consider only phi, bitcast, addrspacecast, and
305// getelementptr operators.
306static bool isAddressExpression(const Value &V, const DataLayout &DL,
307 const TargetTransformInfo *TTI) {
308 const Operator *Op = dyn_cast<Operator>(&V);
309 if (!Op)
310 return false;
311
312 switch (Op->getOpcode()) {
313 case Instruction::PHI:
314 assert(Op->getType()->isPtrOrPtrVectorTy());
315 return true;
316 case Instruction::BitCast:
317 case Instruction::AddrSpaceCast:
318 case Instruction::GetElementPtr:
319 return true;
320 case Instruction::Select:
321 return Op->getType()->isPtrOrPtrVectorTy();
322 case Instruction::Call: {
323 const IntrinsicInst *II = dyn_cast<IntrinsicInst>(&V);
324 return II && II->getIntrinsicID() == Intrinsic::ptrmask;
325 }
326 case Instruction::IntToPtr:
327 return isNoopPtrIntCastPair(Op, DL, TTI);
328 default:
329 // That value is an address expression if it has an assumed address space.
331 }
332}
333
334// Returns the pointer operands of V.
335//
336// Precondition: V is an address expression.
339 const TargetTransformInfo *TTI) {
340 const Operator &Op = cast<Operator>(V);
341 switch (Op.getOpcode()) {
342 case Instruction::PHI: {
343 auto IncomingValues = cast<PHINode>(Op).incoming_values();
344 return {IncomingValues.begin(), IncomingValues.end()};
345 }
346 case Instruction::BitCast:
347 case Instruction::AddrSpaceCast:
348 case Instruction::GetElementPtr:
349 return {Op.getOperand(0)};
350 case Instruction::Select:
351 return {Op.getOperand(1), Op.getOperand(2)};
352 case Instruction::Call: {
353 const IntrinsicInst &II = cast<IntrinsicInst>(Op);
354 assert(II.getIntrinsicID() == Intrinsic::ptrmask &&
355 "unexpected intrinsic call");
356 return {II.getArgOperand(0)};
357 }
358 case Instruction::IntToPtr: {
360 auto *P2I = cast<Operator>(Op.getOperand(0));
361 return {P2I->getOperand(0)};
362 }
363 default:
364 llvm_unreachable("Unexpected instruction type.");
365 }
366}
367
368bool InferAddressSpacesImpl::rewriteIntrinsicOperands(IntrinsicInst *II,
369 Value *OldV,
370 Value *NewV) const {
371 Module *M = II->getParent()->getParent()->getParent();
372
373 switch (II->getIntrinsicID()) {
374 case Intrinsic::objectsize: {
375 Type *DestTy = II->getType();
376 Type *SrcTy = NewV->getType();
377 Function *NewDecl =
378 Intrinsic::getDeclaration(M, II->getIntrinsicID(), {DestTy, SrcTy});
379 II->setArgOperand(0, NewV);
380 II->setCalledFunction(NewDecl);
381 return true;
382 }
383 case Intrinsic::ptrmask:
384 // This is handled as an address expression, not as a use memory operation.
385 return false;
386 case Intrinsic::masked_gather: {
387 Type *RetTy = II->getType();
388 Type *NewPtrTy = NewV->getType();
389 Function *NewDecl =
390 Intrinsic::getDeclaration(M, II->getIntrinsicID(), {RetTy, NewPtrTy});
391 II->setArgOperand(0, NewV);
392 II->setCalledFunction(NewDecl);
393 return true;
394 }
395 case Intrinsic::masked_scatter: {
396 Type *ValueTy = II->getOperand(0)->getType();
397 Type *NewPtrTy = NewV->getType();
398 Function *NewDecl =
399 Intrinsic::getDeclaration(M, II->getIntrinsicID(), {ValueTy, NewPtrTy});
400 II->setArgOperand(1, NewV);
401 II->setCalledFunction(NewDecl);
402 return true;
403 }
404 default: {
405 Value *Rewrite = TTI->rewriteIntrinsicWithAddressSpace(II, OldV, NewV);
406 if (!Rewrite)
407 return false;
408 if (Rewrite != II)
409 II->replaceAllUsesWith(Rewrite);
410 return true;
411 }
412 }
413}
414
415void InferAddressSpacesImpl::collectRewritableIntrinsicOperands(
416 IntrinsicInst *II, PostorderStackTy &PostorderStack,
417 DenseSet<Value *> &Visited) const {
418 auto IID = II->getIntrinsicID();
419 switch (IID) {
420 case Intrinsic::ptrmask:
421 case Intrinsic::objectsize:
422 appendsFlatAddressExpressionToPostorderStack(II->getArgOperand(0),
423 PostorderStack, Visited);
424 break;
425 case Intrinsic::masked_gather:
426 appendsFlatAddressExpressionToPostorderStack(II->getArgOperand(0),
427 PostorderStack, Visited);
428 break;
429 case Intrinsic::masked_scatter:
430 appendsFlatAddressExpressionToPostorderStack(II->getArgOperand(1),
431 PostorderStack, Visited);
432 break;
433 default:
434 SmallVector<int, 2> OpIndexes;
435 if (TTI->collectFlatAddressOperands(OpIndexes, IID)) {
436 for (int Idx : OpIndexes) {
437 appendsFlatAddressExpressionToPostorderStack(II->getArgOperand(Idx),
438 PostorderStack, Visited);
439 }
440 }
441 break;
442 }
443}
444
445// Returns all flat address expressions in function F. The elements are
446// If V is an unvisited flat address expression, appends V to PostorderStack
447// and marks it as visited.
448void InferAddressSpacesImpl::appendsFlatAddressExpressionToPostorderStack(
449 Value *V, PostorderStackTy &PostorderStack,
450 DenseSet<Value *> &Visited) const {
451 assert(V->getType()->isPtrOrPtrVectorTy());
452
453 // Generic addressing expressions may be hidden in nested constant
454 // expressions.
455 if (ConstantExpr *CE = dyn_cast<ConstantExpr>(V)) {
456 // TODO: Look in non-address parts, like icmp operands.
457 if (isAddressExpression(*CE, *DL, TTI) && Visited.insert(CE).second)
458 PostorderStack.emplace_back(CE, false);
459
460 return;
461 }
462
463 if (V->getType()->getPointerAddressSpace() == FlatAddrSpace &&
464 isAddressExpression(*V, *DL, TTI)) {
465 if (Visited.insert(V).second) {
466 PostorderStack.emplace_back(V, false);
467
468 Operator *Op = cast<Operator>(V);
469 for (unsigned I = 0, E = Op->getNumOperands(); I != E; ++I) {
470 if (ConstantExpr *CE = dyn_cast<ConstantExpr>(Op->getOperand(I))) {
471 if (isAddressExpression(*CE, *DL, TTI) && Visited.insert(CE).second)
472 PostorderStack.emplace_back(CE, false);
473 }
474 }
475 }
476 }
477}
478
479// Returns all flat address expressions in function F. The elements are ordered
480// ordered in postorder.
481std::vector<WeakTrackingVH>
482InferAddressSpacesImpl::collectFlatAddressExpressions(Function &F) const {
483 // This function implements a non-recursive postorder traversal of a partial
484 // use-def graph of function F.
485 PostorderStackTy PostorderStack;
486 // The set of visited expressions.
487 DenseSet<Value *> Visited;
488
489 auto PushPtrOperand = [&](Value *Ptr) {
490 appendsFlatAddressExpressionToPostorderStack(Ptr, PostorderStack, Visited);
491 };
492
493 // Look at operations that may be interesting accelerate by moving to a known
494 // address space. We aim at generating after loads and stores, but pure
495 // addressing calculations may also be faster.
496 for (Instruction &I : instructions(F)) {
497 if (auto *GEP = dyn_cast<GetElementPtrInst>(&I)) {
498 PushPtrOperand(GEP->getPointerOperand());
499 } else if (auto *LI = dyn_cast<LoadInst>(&I))
500 PushPtrOperand(LI->getPointerOperand());
501 else if (auto *SI = dyn_cast<StoreInst>(&I))
502 PushPtrOperand(SI->getPointerOperand());
503 else if (auto *RMW = dyn_cast<AtomicRMWInst>(&I))
504 PushPtrOperand(RMW->getPointerOperand());
505 else if (auto *CmpX = dyn_cast<AtomicCmpXchgInst>(&I))
506 PushPtrOperand(CmpX->getPointerOperand());
507 else if (auto *MI = dyn_cast<MemIntrinsic>(&I)) {
508 // For memset/memcpy/memmove, any pointer operand can be replaced.
509 PushPtrOperand(MI->getRawDest());
510
511 // Handle 2nd operand for memcpy/memmove.
512 if (auto *MTI = dyn_cast<MemTransferInst>(MI))
513 PushPtrOperand(MTI->getRawSource());
514 } else if (auto *II = dyn_cast<IntrinsicInst>(&I))
515 collectRewritableIntrinsicOperands(II, PostorderStack, Visited);
516 else if (ICmpInst *Cmp = dyn_cast<ICmpInst>(&I)) {
517 if (Cmp->getOperand(0)->getType()->isPtrOrPtrVectorTy()) {
518 PushPtrOperand(Cmp->getOperand(0));
519 PushPtrOperand(Cmp->getOperand(1));
520 }
521 } else if (auto *ASC = dyn_cast<AddrSpaceCastInst>(&I)) {
522 PushPtrOperand(ASC->getPointerOperand());
523 } else if (auto *I2P = dyn_cast<IntToPtrInst>(&I)) {
524 if (isNoopPtrIntCastPair(cast<Operator>(I2P), *DL, TTI))
525 PushPtrOperand(cast<Operator>(I2P->getOperand(0))->getOperand(0));
526 }
527 }
528
529 std::vector<WeakTrackingVH> Postorder; // The resultant postorder.
530 while (!PostorderStack.empty()) {
531 Value *TopVal = PostorderStack.back().getPointer();
532 // If the operands of the expression on the top are already explored,
533 // adds that expression to the resultant postorder.
534 if (PostorderStack.back().getInt()) {
535 if (TopVal->getType()->getPointerAddressSpace() == FlatAddrSpace)
536 Postorder.push_back(TopVal);
537 PostorderStack.pop_back();
538 continue;
539 }
540 // Otherwise, adds its operands to the stack and explores them.
541 PostorderStack.back().setInt(true);
542 // Skip values with an assumed address space.
544 for (Value *PtrOperand : getPointerOperands(*TopVal, *DL, TTI)) {
545 appendsFlatAddressExpressionToPostorderStack(PtrOperand, PostorderStack,
546 Visited);
547 }
548 }
549 }
550 return Postorder;
551}
552
553// A helper function for cloneInstructionWithNewAddressSpace. Returns the clone
554// of OperandUse.get() in the new address space. If the clone is not ready yet,
555// returns poison in the new address space as a placeholder.
557 const Use &OperandUse, unsigned NewAddrSpace,
558 const ValueToValueMapTy &ValueWithNewAddrSpace,
559 const PredicatedAddrSpaceMapTy &PredicatedAS,
560 SmallVectorImpl<const Use *> *PoisonUsesToFix) {
561 Value *Operand = OperandUse.get();
562
563 Type *NewPtrTy = getPtrOrVecOfPtrsWithNewAS(Operand->getType(), NewAddrSpace);
564
565 if (Constant *C = dyn_cast<Constant>(Operand))
566 return ConstantExpr::getAddrSpaceCast(C, NewPtrTy);
567
568 if (Value *NewOperand = ValueWithNewAddrSpace.lookup(Operand))
569 return NewOperand;
570
571 Instruction *Inst = cast<Instruction>(OperandUse.getUser());
572 auto I = PredicatedAS.find(std::make_pair(Inst, Operand));
573 if (I != PredicatedAS.end()) {
574 // Insert an addrspacecast on that operand before the user.
575 unsigned NewAS = I->second;
576 Type *NewPtrTy = getPtrOrVecOfPtrsWithNewAS(Operand->getType(), NewAS);
577 auto *NewI = new AddrSpaceCastInst(Operand, NewPtrTy);
578 NewI->insertBefore(Inst);
579 NewI->setDebugLoc(Inst->getDebugLoc());
580 return NewI;
581 }
582
583 PoisonUsesToFix->push_back(&OperandUse);
584 return PoisonValue::get(NewPtrTy);
585}
586
587// Returns a clone of `I` with its operands converted to those specified in
588// ValueWithNewAddrSpace. Due to potential cycles in the data flow graph, an
589// operand whose address space needs to be modified might not exist in
590// ValueWithNewAddrSpace. In that case, uses poison as a placeholder operand and
591// adds that operand use to PoisonUsesToFix so that caller can fix them later.
592//
593// Note that we do not necessarily clone `I`, e.g., if it is an addrspacecast
594// from a pointer whose type already matches. Therefore, this function returns a
595// Value* instead of an Instruction*.
596//
597// This may also return nullptr in the case the instruction could not be
598// rewritten.
599Value *InferAddressSpacesImpl::cloneInstructionWithNewAddressSpace(
600 Instruction *I, unsigned NewAddrSpace,
601 const ValueToValueMapTy &ValueWithNewAddrSpace,
602 const PredicatedAddrSpaceMapTy &PredicatedAS,
603 SmallVectorImpl<const Use *> *PoisonUsesToFix) const {
604 Type *NewPtrType = getPtrOrVecOfPtrsWithNewAS(I->getType(), NewAddrSpace);
605
606 if (I->getOpcode() == Instruction::AddrSpaceCast) {
607 Value *Src = I->getOperand(0);
608 // Because `I` is flat, the source address space must be specific.
609 // Therefore, the inferred address space must be the source space, according
610 // to our algorithm.
611 assert(Src->getType()->getPointerAddressSpace() == NewAddrSpace);
612 if (Src->getType() != NewPtrType)
613 return new BitCastInst(Src, NewPtrType);
614 return Src;
615 }
616
617 if (IntrinsicInst *II = dyn_cast<IntrinsicInst>(I)) {
618 // Technically the intrinsic ID is a pointer typed argument, so specially
619 // handle calls early.
620 assert(II->getIntrinsicID() == Intrinsic::ptrmask);
622 II->getArgOperandUse(0), NewAddrSpace, ValueWithNewAddrSpace,
623 PredicatedAS, PoisonUsesToFix);
624 Value *Rewrite =
626 if (Rewrite) {
627 assert(Rewrite != II && "cannot modify this pointer operation in place");
628 return Rewrite;
629 }
630
631 return nullptr;
632 }
633
634 unsigned AS = TTI->getAssumedAddrSpace(I);
635 if (AS != UninitializedAddressSpace) {
636 // For the assumed address space, insert an `addrspacecast` to make that
637 // explicit.
638 Type *NewPtrTy = getPtrOrVecOfPtrsWithNewAS(I->getType(), AS);
639 auto *NewI = new AddrSpaceCastInst(I, NewPtrTy);
640 NewI->insertAfter(I);
641 return NewI;
642 }
643
644 // Computes the converted pointer operands.
645 SmallVector<Value *, 4> NewPointerOperands;
646 for (const Use &OperandUse : I->operands()) {
647 if (!OperandUse.get()->getType()->isPtrOrPtrVectorTy())
648 NewPointerOperands.push_back(nullptr);
649 else
651 OperandUse, NewAddrSpace, ValueWithNewAddrSpace, PredicatedAS,
652 PoisonUsesToFix));
653 }
654
655 switch (I->getOpcode()) {
656 case Instruction::BitCast:
657 return new BitCastInst(NewPointerOperands[0], NewPtrType);
658 case Instruction::PHI: {
659 assert(I->getType()->isPtrOrPtrVectorTy());
660 PHINode *PHI = cast<PHINode>(I);
661 PHINode *NewPHI = PHINode::Create(NewPtrType, PHI->getNumIncomingValues());
662 for (unsigned Index = 0; Index < PHI->getNumIncomingValues(); ++Index) {
663 unsigned OperandNo = PHINode::getOperandNumForIncomingValue(Index);
664 NewPHI->addIncoming(NewPointerOperands[OperandNo],
665 PHI->getIncomingBlock(Index));
666 }
667 return NewPHI;
668 }
669 case Instruction::GetElementPtr: {
670 GetElementPtrInst *GEP = cast<GetElementPtrInst>(I);
672 GEP->getSourceElementType(), NewPointerOperands[0],
673 SmallVector<Value *, 4>(GEP->indices()));
674 NewGEP->setIsInBounds(GEP->isInBounds());
675 return NewGEP;
676 }
677 case Instruction::Select:
678 assert(I->getType()->isPtrOrPtrVectorTy());
679 return SelectInst::Create(I->getOperand(0), NewPointerOperands[1],
680 NewPointerOperands[2], "", nullptr, I);
681 case Instruction::IntToPtr: {
682 assert(isNoopPtrIntCastPair(cast<Operator>(I), *DL, TTI));
683 Value *Src = cast<Operator>(I->getOperand(0))->getOperand(0);
684 if (Src->getType() == NewPtrType)
685 return Src;
686
687 // If we had a no-op inttoptr/ptrtoint pair, we may still have inferred a
688 // source address space from a generic pointer source need to insert a cast
689 // back.
690 return CastInst::CreatePointerBitCastOrAddrSpaceCast(Src, NewPtrType);
691 }
692 default:
693 llvm_unreachable("Unexpected opcode");
694 }
695}
696
697// Similar to cloneInstructionWithNewAddressSpace, returns a clone of the
698// constant expression `CE` with its operands replaced as specified in
699// ValueWithNewAddrSpace.
701 ConstantExpr *CE, unsigned NewAddrSpace,
702 const ValueToValueMapTy &ValueWithNewAddrSpace, const DataLayout *DL,
703 const TargetTransformInfo *TTI) {
704 Type *TargetType =
705 CE->getType()->isPtrOrPtrVectorTy()
706 ? getPtrOrVecOfPtrsWithNewAS(CE->getType(), NewAddrSpace)
707 : CE->getType();
708
709 if (CE->getOpcode() == Instruction::AddrSpaceCast) {
710 // Because CE is flat, the source address space must be specific.
711 // Therefore, the inferred address space must be the source space according
712 // to our algorithm.
713 assert(CE->getOperand(0)->getType()->getPointerAddressSpace() ==
714 NewAddrSpace);
715 return ConstantExpr::getBitCast(CE->getOperand(0), TargetType);
716 }
717
718 if (CE->getOpcode() == Instruction::BitCast) {
719 if (Value *NewOperand = ValueWithNewAddrSpace.lookup(CE->getOperand(0)))
720 return ConstantExpr::getBitCast(cast<Constant>(NewOperand), TargetType);
721 return ConstantExpr::getAddrSpaceCast(CE, TargetType);
722 }
723
724 if (CE->getOpcode() == Instruction::IntToPtr) {
725 assert(isNoopPtrIntCastPair(cast<Operator>(CE), *DL, TTI));
726 Constant *Src = cast<ConstantExpr>(CE->getOperand(0))->getOperand(0);
727 assert(Src->getType()->getPointerAddressSpace() == NewAddrSpace);
728 return ConstantExpr::getBitCast(Src, TargetType);
729 }
730
731 // Computes the operands of the new constant expression.
732 bool IsNew = false;
733 SmallVector<Constant *, 4> NewOperands;
734 for (unsigned Index = 0; Index < CE->getNumOperands(); ++Index) {
735 Constant *Operand = CE->getOperand(Index);
736 // If the address space of `Operand` needs to be modified, the new operand
737 // with the new address space should already be in ValueWithNewAddrSpace
738 // because (1) the constant expressions we consider (i.e. addrspacecast,
739 // bitcast, and getelementptr) do not incur cycles in the data flow graph
740 // and (2) this function is called on constant expressions in postorder.
741 if (Value *NewOperand = ValueWithNewAddrSpace.lookup(Operand)) {
742 IsNew = true;
743 NewOperands.push_back(cast<Constant>(NewOperand));
744 continue;
745 }
746 if (auto *CExpr = dyn_cast<ConstantExpr>(Operand))
748 CExpr, NewAddrSpace, ValueWithNewAddrSpace, DL, TTI)) {
749 IsNew = true;
750 NewOperands.push_back(cast<Constant>(NewOperand));
751 continue;
752 }
753 // Otherwise, reuses the old operand.
754 NewOperands.push_back(Operand);
755 }
756
757 // If !IsNew, we will replace the Value with itself. However, replaced values
758 // are assumed to wrapped in an addrspacecast cast later so drop it now.
759 if (!IsNew)
760 return nullptr;
761
762 if (CE->getOpcode() == Instruction::GetElementPtr) {
763 // Needs to specify the source type while constructing a getelementptr
764 // constant expression.
765 return CE->getWithOperands(NewOperands, TargetType, /*OnlyIfReduced=*/false,
766 cast<GEPOperator>(CE)->getSourceElementType());
767 }
768
769 return CE->getWithOperands(NewOperands, TargetType);
770}
771
772// Returns a clone of the value `V`, with its operands replaced as specified in
773// ValueWithNewAddrSpace. This function is called on every flat address
774// expression whose address space needs to be modified, in postorder.
775//
776// See cloneInstructionWithNewAddressSpace for the meaning of PoisonUsesToFix.
777Value *InferAddressSpacesImpl::cloneValueWithNewAddressSpace(
778 Value *V, unsigned NewAddrSpace,
779 const ValueToValueMapTy &ValueWithNewAddrSpace,
780 const PredicatedAddrSpaceMapTy &PredicatedAS,
781 SmallVectorImpl<const Use *> *PoisonUsesToFix) const {
782 // All values in Postorder are flat address expressions.
783 assert(V->getType()->getPointerAddressSpace() == FlatAddrSpace &&
784 isAddressExpression(*V, *DL, TTI));
785
786 if (Instruction *I = dyn_cast<Instruction>(V)) {
787 Value *NewV = cloneInstructionWithNewAddressSpace(
788 I, NewAddrSpace, ValueWithNewAddrSpace, PredicatedAS, PoisonUsesToFix);
789 if (Instruction *NewI = dyn_cast_or_null<Instruction>(NewV)) {
790 if (NewI->getParent() == nullptr) {
791 NewI->insertBefore(I);
792 NewI->takeName(I);
793 NewI->setDebugLoc(I->getDebugLoc());
794 }
795 }
796 return NewV;
797 }
798
800 cast<ConstantExpr>(V), NewAddrSpace, ValueWithNewAddrSpace, DL, TTI);
801}
802
803// Defines the join operation on the address space lattice (see the file header
804// comments).
805unsigned InferAddressSpacesImpl::joinAddressSpaces(unsigned AS1,
806 unsigned AS2) const {
807 if (AS1 == FlatAddrSpace || AS2 == FlatAddrSpace)
808 return FlatAddrSpace;
809
810 if (AS1 == UninitializedAddressSpace)
811 return AS2;
812 if (AS2 == UninitializedAddressSpace)
813 return AS1;
814
815 // The join of two different specific address spaces is flat.
816 return (AS1 == AS2) ? AS1 : FlatAddrSpace;
817}
818
819bool InferAddressSpacesImpl::run(Function &F) {
820 DL = &F.getParent()->getDataLayout();
821
823 FlatAddrSpace = 0;
824
825 if (FlatAddrSpace == UninitializedAddressSpace) {
826 FlatAddrSpace = TTI->getFlatAddressSpace();
827 if (FlatAddrSpace == UninitializedAddressSpace)
828 return false;
829 }
830
831 // Collects all flat address expressions in postorder.
832 std::vector<WeakTrackingVH> Postorder = collectFlatAddressExpressions(F);
833
834 // Runs a data-flow analysis to refine the address spaces of every expression
835 // in Postorder.
836 ValueToAddrSpaceMapTy InferredAddrSpace;
837 PredicatedAddrSpaceMapTy PredicatedAS;
838 inferAddressSpaces(Postorder, InferredAddrSpace, PredicatedAS);
839
840 // Changes the address spaces of the flat address expressions who are inferred
841 // to point to a specific address space.
842 return rewriteWithNewAddressSpaces(Postorder, InferredAddrSpace, PredicatedAS,
843 &F);
844}
845
846// Constants need to be tracked through RAUW to handle cases with nested
847// constant expressions, so wrap values in WeakTrackingVH.
848void InferAddressSpacesImpl::inferAddressSpaces(
849 ArrayRef<WeakTrackingVH> Postorder,
850 ValueToAddrSpaceMapTy &InferredAddrSpace,
851 PredicatedAddrSpaceMapTy &PredicatedAS) const {
852 SetVector<Value *> Worklist(Postorder.begin(), Postorder.end());
853 // Initially, all expressions are in the uninitialized address space.
854 for (Value *V : Postorder)
855 InferredAddrSpace[V] = UninitializedAddressSpace;
856
857 while (!Worklist.empty()) {
858 Value *V = Worklist.pop_back_val();
859
860 // Try to update the address space of the stack top according to the
861 // address spaces of its operands.
862 if (!updateAddressSpace(*V, InferredAddrSpace, PredicatedAS))
863 continue;
864
865 for (Value *User : V->users()) {
866 // Skip if User is already in the worklist.
867 if (Worklist.count(User))
868 continue;
869
870 auto Pos = InferredAddrSpace.find(User);
871 // Our algorithm only updates the address spaces of flat address
872 // expressions, which are those in InferredAddrSpace.
873 if (Pos == InferredAddrSpace.end())
874 continue;
875
876 // Function updateAddressSpace moves the address space down a lattice
877 // path. Therefore, nothing to do if User is already inferred as flat (the
878 // bottom element in the lattice).
879 if (Pos->second == FlatAddrSpace)
880 continue;
881
882 Worklist.insert(User);
883 }
884 }
885}
886
887unsigned InferAddressSpacesImpl::getPredicatedAddrSpace(const Value &V,
888 Value *Opnd) const {
889 const Instruction *I = dyn_cast<Instruction>(&V);
890 if (!I)
892
893 Opnd = Opnd->stripInBoundsOffsets();
894 for (auto &AssumeVH : AC.assumptionsFor(Opnd)) {
895 if (!AssumeVH)
896 continue;
897 CallInst *CI = cast<CallInst>(AssumeVH);
898 if (!isValidAssumeForContext(CI, I, DT))
899 continue;
900
901 const Value *Ptr;
902 unsigned AS;
903 std::tie(Ptr, AS) = TTI->getPredicatedAddrSpace(CI->getArgOperand(0));
904 if (Ptr)
905 return AS;
906 }
907
909}
910
911bool InferAddressSpacesImpl::updateAddressSpace(
912 const Value &V, ValueToAddrSpaceMapTy &InferredAddrSpace,
913 PredicatedAddrSpaceMapTy &PredicatedAS) const {
914 assert(InferredAddrSpace.count(&V));
915
916 LLVM_DEBUG(dbgs() << "Updating the address space of\n " << V << '\n');
917
918 // The new inferred address space equals the join of the address spaces
919 // of all its pointer operands.
920 unsigned NewAS = UninitializedAddressSpace;
921
922 const Operator &Op = cast<Operator>(V);
923 if (Op.getOpcode() == Instruction::Select) {
924 Value *Src0 = Op.getOperand(1);
925 Value *Src1 = Op.getOperand(2);
926
927 auto I = InferredAddrSpace.find(Src0);
928 unsigned Src0AS = (I != InferredAddrSpace.end())
929 ? I->second
930 : Src0->getType()->getPointerAddressSpace();
931
932 auto J = InferredAddrSpace.find(Src1);
933 unsigned Src1AS = (J != InferredAddrSpace.end())
934 ? J->second
935 : Src1->getType()->getPointerAddressSpace();
936
937 auto *C0 = dyn_cast<Constant>(Src0);
938 auto *C1 = dyn_cast<Constant>(Src1);
939
940 // If one of the inputs is a constant, we may be able to do a constant
941 // addrspacecast of it. Defer inferring the address space until the input
942 // address space is known.
943 if ((C1 && Src0AS == UninitializedAddressSpace) ||
944 (C0 && Src1AS == UninitializedAddressSpace))
945 return false;
946
947 if (C0 && isSafeToCastConstAddrSpace(C0, Src1AS))
948 NewAS = Src1AS;
949 else if (C1 && isSafeToCastConstAddrSpace(C1, Src0AS))
950 NewAS = Src0AS;
951 else
952 NewAS = joinAddressSpaces(Src0AS, Src1AS);
953 } else {
954 unsigned AS = TTI->getAssumedAddrSpace(&V);
955 if (AS != UninitializedAddressSpace) {
956 // Use the assumed address space directly.
957 NewAS = AS;
958 } else {
959 // Otherwise, infer the address space from its pointer operands.
960 for (Value *PtrOperand : getPointerOperands(V, *DL, TTI)) {
961 auto I = InferredAddrSpace.find(PtrOperand);
962 unsigned OperandAS;
963 if (I == InferredAddrSpace.end()) {
964 OperandAS = PtrOperand->getType()->getPointerAddressSpace();
965 if (OperandAS == FlatAddrSpace) {
966 // Check AC for assumption dominating V.
967 unsigned AS = getPredicatedAddrSpace(V, PtrOperand);
968 if (AS != UninitializedAddressSpace) {
970 << " deduce operand AS from the predicate addrspace "
971 << AS << '\n');
972 OperandAS = AS;
973 // Record this use with the predicated AS.
974 PredicatedAS[std::make_pair(&V, PtrOperand)] = OperandAS;
975 }
976 }
977 } else
978 OperandAS = I->second;
979
980 // join(flat, *) = flat. So we can break if NewAS is already flat.
981 NewAS = joinAddressSpaces(NewAS, OperandAS);
982 if (NewAS == FlatAddrSpace)
983 break;
984 }
985 }
986 }
987
988 unsigned OldAS = InferredAddrSpace.lookup(&V);
989 assert(OldAS != FlatAddrSpace);
990 if (OldAS == NewAS)
991 return false;
992
993 // If any updates are made, grabs its users to the worklist because
994 // their address spaces can also be possibly updated.
995 LLVM_DEBUG(dbgs() << " to " << NewAS << '\n');
996 InferredAddrSpace[&V] = NewAS;
997 return true;
998}
999
1000/// \p returns true if \p U is the pointer operand of a memory instruction with
1001/// a single pointer operand that can have its address space changed by simply
1002/// mutating the use to a new value. If the memory instruction is volatile,
1003/// return true only if the target allows the memory instruction to be volatile
1004/// in the new address space.
1006 Use &U, unsigned AddrSpace) {
1007 User *Inst = U.getUser();
1008 unsigned OpNo = U.getOperandNo();
1009 bool VolatileIsAllowed = false;
1010 if (auto *I = dyn_cast<Instruction>(Inst))
1011 VolatileIsAllowed = TTI.hasVolatileVariant(I, AddrSpace);
1012
1013 if (auto *LI = dyn_cast<LoadInst>(Inst))
1014 return OpNo == LoadInst::getPointerOperandIndex() &&
1015 (VolatileIsAllowed || !LI->isVolatile());
1016
1017 if (auto *SI = dyn_cast<StoreInst>(Inst))
1018 return OpNo == StoreInst::getPointerOperandIndex() &&
1019 (VolatileIsAllowed || !SI->isVolatile());
1020
1021 if (auto *RMW = dyn_cast<AtomicRMWInst>(Inst))
1022 return OpNo == AtomicRMWInst::getPointerOperandIndex() &&
1023 (VolatileIsAllowed || !RMW->isVolatile());
1024
1025 if (auto *CmpX = dyn_cast<AtomicCmpXchgInst>(Inst))
1027 (VolatileIsAllowed || !CmpX->isVolatile());
1028
1029 return false;
1030}
1031
1032/// Update memory intrinsic uses that require more complex processing than
1033/// simple memory instructions. These require re-mangling and may have multiple
1034/// pointer operands.
1036 Value *NewV) {
1037 IRBuilder<> B(MI);
1038 MDNode *TBAA = MI->getMetadata(LLVMContext::MD_tbaa);
1039 MDNode *ScopeMD = MI->getMetadata(LLVMContext::MD_alias_scope);
1040 MDNode *NoAliasMD = MI->getMetadata(LLVMContext::MD_noalias);
1041
1042 if (auto *MSI = dyn_cast<MemSetInst>(MI)) {
1043 B.CreateMemSet(NewV, MSI->getValue(), MSI->getLength(), MSI->getDestAlign(),
1044 false, // isVolatile
1045 TBAA, ScopeMD, NoAliasMD);
1046 } else if (auto *MTI = dyn_cast<MemTransferInst>(MI)) {
1047 Value *Src = MTI->getRawSource();
1048 Value *Dest = MTI->getRawDest();
1049
1050 // Be careful in case this is a self-to-self copy.
1051 if (Src == OldV)
1052 Src = NewV;
1053
1054 if (Dest == OldV)
1055 Dest = NewV;
1056
1057 if (isa<MemCpyInlineInst>(MTI)) {
1058 MDNode *TBAAStruct = MTI->getMetadata(LLVMContext::MD_tbaa_struct);
1059 B.CreateMemCpyInline(Dest, MTI->getDestAlign(), Src,
1060 MTI->getSourceAlign(), MTI->getLength(),
1061 false, // isVolatile
1062 TBAA, TBAAStruct, ScopeMD, NoAliasMD);
1063 } else if (isa<MemCpyInst>(MTI)) {
1064 MDNode *TBAAStruct = MTI->getMetadata(LLVMContext::MD_tbaa_struct);
1065 B.CreateMemCpy(Dest, MTI->getDestAlign(), Src, MTI->getSourceAlign(),
1066 MTI->getLength(),
1067 false, // isVolatile
1068 TBAA, TBAAStruct, ScopeMD, NoAliasMD);
1069 } else {
1070 assert(isa<MemMoveInst>(MTI));
1071 B.CreateMemMove(Dest, MTI->getDestAlign(), Src, MTI->getSourceAlign(),
1072 MTI->getLength(),
1073 false, // isVolatile
1074 TBAA, ScopeMD, NoAliasMD);
1075 }
1076 } else
1077 llvm_unreachable("unhandled MemIntrinsic");
1078
1079 MI->eraseFromParent();
1080 return true;
1081}
1082
1083// \p returns true if it is OK to change the address space of constant \p C with
1084// a ConstantExpr addrspacecast.
1085bool InferAddressSpacesImpl::isSafeToCastConstAddrSpace(Constant *C,
1086 unsigned NewAS) const {
1088
1089 unsigned SrcAS = C->getType()->getPointerAddressSpace();
1090 if (SrcAS == NewAS || isa<UndefValue>(C))
1091 return true;
1092
1093 // Prevent illegal casts between different non-flat address spaces.
1094 if (SrcAS != FlatAddrSpace && NewAS != FlatAddrSpace)
1095 return false;
1096
1097 if (isa<ConstantPointerNull>(C))
1098 return true;
1099
1100 if (auto *Op = dyn_cast<Operator>(C)) {
1101 // If we already have a constant addrspacecast, it should be safe to cast it
1102 // off.
1103 if (Op->getOpcode() == Instruction::AddrSpaceCast)
1104 return isSafeToCastConstAddrSpace(cast<Constant>(Op->getOperand(0)),
1105 NewAS);
1106
1107 if (Op->getOpcode() == Instruction::IntToPtr &&
1108 Op->getType()->getPointerAddressSpace() == FlatAddrSpace)
1109 return true;
1110 }
1111
1112 return false;
1113}
1114
1117 User *CurUser = I->getUser();
1118 ++I;
1119
1120 while (I != End && I->getUser() == CurUser)
1121 ++I;
1122
1123 return I;
1124}
1125
1126bool InferAddressSpacesImpl::rewriteWithNewAddressSpaces(
1127 ArrayRef<WeakTrackingVH> Postorder,
1128 const ValueToAddrSpaceMapTy &InferredAddrSpace,
1129 const PredicatedAddrSpaceMapTy &PredicatedAS, Function *F) const {
1130 // For each address expression to be modified, creates a clone of it with its
1131 // pointer operands converted to the new address space. Since the pointer
1132 // operands are converted, the clone is naturally in the new address space by
1133 // construction.
1134 ValueToValueMapTy ValueWithNewAddrSpace;
1135 SmallVector<const Use *, 32> PoisonUsesToFix;
1136 for (Value *V : Postorder) {
1137 unsigned NewAddrSpace = InferredAddrSpace.lookup(V);
1138
1139 // In some degenerate cases (e.g. invalid IR in unreachable code), we may
1140 // not even infer the value to have its original address space.
1141 if (NewAddrSpace == UninitializedAddressSpace)
1142 continue;
1143
1144 if (V->getType()->getPointerAddressSpace() != NewAddrSpace) {
1145 Value *New =
1146 cloneValueWithNewAddressSpace(V, NewAddrSpace, ValueWithNewAddrSpace,
1147 PredicatedAS, &PoisonUsesToFix);
1148 if (New)
1149 ValueWithNewAddrSpace[V] = New;
1150 }
1151 }
1152
1153 if (ValueWithNewAddrSpace.empty())
1154 return false;
1155
1156 // Fixes all the poison uses generated by cloneInstructionWithNewAddressSpace.
1157 for (const Use *PoisonUse : PoisonUsesToFix) {
1158 User *V = PoisonUse->getUser();
1159 User *NewV = cast_or_null<User>(ValueWithNewAddrSpace.lookup(V));
1160 if (!NewV)
1161 continue;
1162
1163 unsigned OperandNo = PoisonUse->getOperandNo();
1164 assert(isa<PoisonValue>(NewV->getOperand(OperandNo)));
1165 NewV->setOperand(OperandNo, ValueWithNewAddrSpace.lookup(PoisonUse->get()));
1166 }
1167
1168 SmallVector<Instruction *, 16> DeadInstructions;
1169
1170 // Replaces the uses of the old address expressions with the new ones.
1171 for (const WeakTrackingVH &WVH : Postorder) {
1172 assert(WVH && "value was unexpectedly deleted");
1173 Value *V = WVH;
1174 Value *NewV = ValueWithNewAddrSpace.lookup(V);
1175 if (NewV == nullptr)
1176 continue;
1177
1178 LLVM_DEBUG(dbgs() << "Replacing the uses of " << *V << "\n with\n "
1179 << *NewV << '\n');
1180
1181 if (Constant *C = dyn_cast<Constant>(V)) {
1182 Constant *Replace =
1183 ConstantExpr::getAddrSpaceCast(cast<Constant>(NewV), C->getType());
1184 if (C != Replace) {
1185 LLVM_DEBUG(dbgs() << "Inserting replacement const cast: " << Replace
1186 << ": " << *Replace << '\n');
1187 C->replaceAllUsesWith(Replace);
1188 V = Replace;
1189 }
1190 }
1191
1192 Value::use_iterator I, E, Next;
1193 for (I = V->use_begin(), E = V->use_end(); I != E;) {
1194 Use &U = *I;
1195
1196 // Some users may see the same pointer operand in multiple operands. Skip
1197 // to the next instruction.
1198 I = skipToNextUser(I, E);
1199
1201 *TTI, U, V->getType()->getPointerAddressSpace())) {
1202 // If V is used as the pointer operand of a compatible memory operation,
1203 // sets the pointer operand to NewV. This replacement does not change
1204 // the element type, so the resultant load/store is still valid.
1205 U.set(NewV);
1206 continue;
1207 }
1208
1209 User *CurUser = U.getUser();
1210 // Skip if the current user is the new value itself.
1211 if (CurUser == NewV)
1212 continue;
1213 // Handle more complex cases like intrinsic that need to be remangled.
1214 if (auto *MI = dyn_cast<MemIntrinsic>(CurUser)) {
1215 if (!MI->isVolatile() && handleMemIntrinsicPtrUse(MI, V, NewV))
1216 continue;
1217 }
1218
1219 if (auto *II = dyn_cast<IntrinsicInst>(CurUser)) {
1220 if (rewriteIntrinsicOperands(II, V, NewV))
1221 continue;
1222 }
1223
1224 if (isa<Instruction>(CurUser)) {
1225 if (ICmpInst *Cmp = dyn_cast<ICmpInst>(CurUser)) {
1226 // If we can infer that both pointers are in the same addrspace,
1227 // transform e.g.
1228 // %cmp = icmp eq float* %p, %q
1229 // into
1230 // %cmp = icmp eq float addrspace(3)* %new_p, %new_q
1231
1232 unsigned NewAS = NewV->getType()->getPointerAddressSpace();
1233 int SrcIdx = U.getOperandNo();
1234 int OtherIdx = (SrcIdx == 0) ? 1 : 0;
1235 Value *OtherSrc = Cmp->getOperand(OtherIdx);
1236
1237 if (Value *OtherNewV = ValueWithNewAddrSpace.lookup(OtherSrc)) {
1238 if (OtherNewV->getType()->getPointerAddressSpace() == NewAS) {
1239 Cmp->setOperand(OtherIdx, OtherNewV);
1240 Cmp->setOperand(SrcIdx, NewV);
1241 continue;
1242 }
1243 }
1244
1245 // Even if the type mismatches, we can cast the constant.
1246 if (auto *KOtherSrc = dyn_cast<Constant>(OtherSrc)) {
1247 if (isSafeToCastConstAddrSpace(KOtherSrc, NewAS)) {
1248 Cmp->setOperand(SrcIdx, NewV);
1249 Cmp->setOperand(OtherIdx, ConstantExpr::getAddrSpaceCast(
1250 KOtherSrc, NewV->getType()));
1251 continue;
1252 }
1253 }
1254 }
1255
1256 if (AddrSpaceCastInst *ASC = dyn_cast<AddrSpaceCastInst>(CurUser)) {
1257 unsigned NewAS = NewV->getType()->getPointerAddressSpace();
1258 if (ASC->getDestAddressSpace() == NewAS) {
1259 ASC->replaceAllUsesWith(NewV);
1260 DeadInstructions.push_back(ASC);
1261 continue;
1262 }
1263 }
1264
1265 // Otherwise, replaces the use with flat(NewV).
1266 if (Instruction *VInst = dyn_cast<Instruction>(V)) {
1267 // Don't create a copy of the original addrspacecast.
1268 if (U == V && isa<AddrSpaceCastInst>(V))
1269 continue;
1270
1271 // Insert the addrspacecast after NewV.
1272 BasicBlock::iterator InsertPos;
1273 if (Instruction *NewVInst = dyn_cast<Instruction>(NewV))
1274 InsertPos = std::next(NewVInst->getIterator());
1275 else
1276 InsertPos = std::next(VInst->getIterator());
1277
1278 while (isa<PHINode>(InsertPos))
1279 ++InsertPos;
1280 U.set(new AddrSpaceCastInst(NewV, V->getType(), "", &*InsertPos));
1281 } else {
1282 U.set(ConstantExpr::getAddrSpaceCast(cast<Constant>(NewV),
1283 V->getType()));
1284 }
1285 }
1286 }
1287
1288 if (V->use_empty()) {
1289 if (Instruction *I = dyn_cast<Instruction>(V))
1290 DeadInstructions.push_back(I);
1291 }
1292 }
1293
1294 for (Instruction *I : DeadInstructions)
1296
1297 return true;
1298}
1299
1300bool InferAddressSpaces::runOnFunction(Function &F) {
1301 if (skipFunction(F))
1302 return false;
1303
1304 auto *DTWP = getAnalysisIfAvailable<DominatorTreeWrapperPass>();
1305 DominatorTree *DT = DTWP ? &DTWP->getDomTree() : nullptr;
1306 return InferAddressSpacesImpl(
1307 getAnalysis<AssumptionCacheTracker>().getAssumptionCache(F), DT,
1308 &getAnalysis<TargetTransformInfoWrapperPass>().getTTI(F),
1309 FlatAddrSpace)
1310 .run(F);
1311}
1312
1314 return new InferAddressSpaces(AddressSpace);
1315}
1316
1318 : FlatAddrSpace(UninitializedAddressSpace) {}
1320 : FlatAddrSpace(AddressSpace) {}
1321
1324 bool Changed =
1325 InferAddressSpacesImpl(AM.getResult<AssumptionAnalysis>(F),
1327 &AM.getResult<TargetIRAnalysis>(F), FlatAddrSpace)
1328 .run(F);
1329 if (Changed) {
1333 return PA;
1334 }
1335 return PreservedAnalyses::all();
1336}
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
Rewrite undef for PHI
static GCRegistry::Add< OcamlGC > B("ocaml", "ocaml 3.10-compatible GC")
static GCRegistry::Add< CoreCLRGC > E("coreclr", "CoreCLR-compatible GC")
This file contains the declarations for the subclasses of Constant, which represent the different fla...
return RetTy
Returns the sub type a function will return at a given Idx Should correspond to the result type of an ExtractValue instruction executed with just that one unsigned Idx
#define LLVM_DEBUG(X)
Definition: Debug.h:101
This file defines the DenseMap class.
This file defines the DenseSet and SmallDenseSet classes.
bool End
Definition: ELF_riscv.cpp:469
Hexagon Common GEP
IRTranslator LLVM IR MI
static cl::opt< bool > AssumeDefaultIsFlatAddressSpace("assume-default-is-flat-addrspace", cl::init(false), cl::ReallyHidden, cl::desc("The default address space is assumed as the flat address space. " "This is mainly for test purpose."))
static bool isNoopPtrIntCastPair(const Operator *I2P, const DataLayout &DL, const TargetTransformInfo *TTI)
static bool isAddressExpression(const Value &V, const DataLayout &DL, const TargetTransformInfo *TTI)
static bool handleMemIntrinsicPtrUse(MemIntrinsic *MI, Value *OldV, Value *NewV)
Update memory intrinsic uses that require more complex processing than simple memory instructions.
Infer address spaces
static SmallVector< Value *, 2 > getPointerOperands(const Value &V, const DataLayout &DL, const TargetTransformInfo *TTI)
static Value * operandWithNewAddressSpaceOrCreatePoison(const Use &OperandUse, unsigned NewAddrSpace, const ValueToValueMapTy &ValueWithNewAddrSpace, const PredicatedAddrSpaceMapTy &PredicatedAS, SmallVectorImpl< const Use * > *PoisonUsesToFix)
static Value * cloneConstantExprWithNewAddressSpace(ConstantExpr *CE, unsigned NewAddrSpace, const ValueToValueMapTy &ValueWithNewAddrSpace, const DataLayout *DL, const TargetTransformInfo *TTI)
static bool isSimplePointerUseValidToReplace(const TargetTransformInfo &TTI, Use &U, unsigned AddrSpace)
returns true if U is the pointer operand of a memory instruction with a single pointer operand that c...
static Value::use_iterator skipToNextUser(Value::use_iterator I, Value::use_iterator End)
Infer address static false Type * getPtrOrVecOfPtrsWithNewAS(Type *Ty, unsigned NewAddrSpace)
#define DEBUG_TYPE
static const unsigned UninitializedAddressSpace
Select target instructions out of generic instructions
#define F(x, y, z)
Definition: MD5.cpp:55
#define I(x, y, z)
Definition: MD5.cpp:58
This header defines various interfaces for pass management in LLVM.
#define INITIALIZE_PASS_DEPENDENCY(depName)
Definition: PassSupport.h:55
#define INITIALIZE_PASS_END(passName, arg, name, cfg, analysis)
Definition: PassSupport.h:59
#define INITIALIZE_PASS_BEGIN(passName, arg, name, cfg, analysis)
Definition: PassSupport.h:52
assert(ImpDefSCC.getReg()==AMDGPU::SCC &&ImpDefSCC.isDef())
This file implements a set that has insertion order iteration characteristics.
This file defines the SmallVector class.
This pass exposes codegen information to IR-level passes.
This defines the Use class.
This class represents a conversion between pointers from one address space to another.
A container for analyses that lazily runs them and caches their results.
Definition: PassManager.h:620
PassT::Result * getCachedResult(IRUnitT &IR) const
Get the cached result of an analysis pass for a given IR unit.
Definition: PassManager.h:793
PassT::Result & getResult(IRUnitT &IR, ExtraArgTs... ExtraArgs)
Get the result of an analysis pass for a given IR unit.
Definition: PassManager.h:774
Represent the analysis usage information of a pass.
AnalysisUsage & addRequired()
AnalysisUsage & addPreserved()
Add the specified Pass class to the set of analyses preserved by this pass.
void setPreservesCFG()
This function should be called by the pass, iff they do not:
Definition: Pass.cpp:269
ArrayRef - Represent a constant reference to an array (0 or more elements consecutively in memory),...
Definition: ArrayRef.h:41
iterator end() const
Definition: ArrayRef.h:154
iterator begin() const
Definition: ArrayRef.h:153
A function analysis which provides an AssumptionCache.
An immutable pass that tracks lazily created AssumptionCache objects.
A cache of @llvm.assume calls within a function.
static unsigned getPointerOperandIndex()
Definition: Instructions.h:645
static unsigned getPointerOperandIndex()
Definition: Instructions.h:879
const Function * getParent() const
Return the enclosing method, or null if none.
Definition: BasicBlock.h:112
InstListType::iterator iterator
Instruction iterators...
Definition: BasicBlock.h:87
This class represents a no-op cast from one type to another.
Represents analyses that only rely on functions' control flow.
Definition: PassManager.h:113
const Use & getArgOperandUse(unsigned i) const
Wrappers for getting the Use of a call argument.
Definition: InstrTypes.h:1368
Value * getArgOperand(unsigned i) const
Definition: InstrTypes.h:1357
void setArgOperand(unsigned i, Value *v)
Definition: InstrTypes.h:1362
void setCalledFunction(Function *Fn)
Sets the function called, including updating the function type.
Definition: InstrTypes.h:1451
This class represents a function call, abstracting a target machine's calling convention.
static CastInst * CreatePointerBitCastOrAddrSpaceCast(Value *S, Type *Ty, const Twine &Name, BasicBlock *InsertAtEnd)
Create a BitCast or an AddrSpaceCast cast instruction.
static bool isNoopCast(Instruction::CastOps Opcode, Type *SrcTy, Type *DstTy, const DataLayout &DL)
A no-op cast is one that can be effected without changing any bits.
A constant value that is initialized with an expression using other constant values.
Definition: Constants.h:997
static Constant * getAddrSpaceCast(Constant *C, Type *Ty, bool OnlyIfReduced=false)
Definition: Constants.cpp:2225
static Constant * getBitCast(Constant *C, Type *Ty, bool OnlyIfReduced=false)
Definition: Constants.cpp:2213
This is an important base class in LLVM.
Definition: Constant.h:41
This class represents an Operation in the Expression.
A parsed version of the target data layout string in and methods for querying it.
Definition: DataLayout.h:110
Implements a dense probed hash-table based set.
Definition: DenseSet.h:271
Analysis pass which computes a DominatorTree.
Definition: Dominators.h:279
Legacy analysis pass which computes a DominatorTree.
Definition: Dominators.h:314
Concrete subclass of DominatorTreeBase that is used to compute a normal dominator tree.
Definition: Dominators.h:166
FunctionPass class - This class is used to implement most global optimizations.
Definition: Pass.h:311
virtual bool runOnFunction(Function &F)=0
runOnFunction - Virtual method overriden by subclasses to do the per-function processing of the pass.
an instruction for type-safe pointer arithmetic to access elements of arrays and structs
Definition: Instructions.h:940
void setIsInBounds(bool b=true)
Set or clear the inbounds flag on this GEP instruction.
static GetElementPtrInst * Create(Type *PointeeType, Value *Ptr, ArrayRef< Value * > IdxList, const Twine &NameStr="", Instruction *InsertBefore=nullptr)
Definition: Instructions.h:966
Module * getParent()
Get the module that this global value is contained inside of...
Definition: GlobalValue.h:652
This instruction compares its operands according to the predicate given to the constructor.
This provides a uniform API for creating instructions and inserting them into a basic block: either a...
Definition: IRBuilder.h:2628
const DebugLoc & getDebugLoc() const
Return the debug location for this node as a DebugLoc.
Definition: Instruction.h:392
const BasicBlock * getParent() const
Definition: Instruction.h:90
A wrapper class for inspecting calls to intrinsic functions.
Definition: IntrinsicInst.h:47
Intrinsic::ID getIntrinsicID() const
Return the intrinsic ID of this intrinsic.
Definition: IntrinsicInst.h:54
static unsigned getPointerOperandIndex()
Definition: Instructions.h:266
Metadata node.
Definition: Metadata.h:950
This is the common base class for memset/memcpy/memmove.
A Module instance is used to store all the information related to an LLVM module.
Definition: Module.h:65
This is a utility class that provides an abstraction for the common functionality between Instruction...
Definition: Operator.h:31
unsigned getOpcode() const
Return the opcode for this Instruction or ConstantExpr.
Definition: Operator.h:41
void addIncoming(Value *V, BasicBlock *BB)
Add an incoming value to the end of the PHI list.
static PHINode * Create(Type *Ty, unsigned NumReservedValues, const Twine &NameStr="", Instruction *InsertBefore=nullptr)
Constructors - NumReservedValues is a hint for the number of incoming edges that this phi node will h...
static unsigned getOperandNumForIncomingValue(unsigned i)
static PassRegistry * getPassRegistry()
getPassRegistry - Access the global registry object, which is automatically initialized at applicatio...
virtual void getAnalysisUsage(AnalysisUsage &) const
getAnalysisUsage - This function should be overriden by passes that need analysis information to do t...
Definition: Pass.cpp:98
static PoisonValue * get(Type *T)
Static factory methods - Return an 'poison' object of the specified type.
Definition: Constants.cpp:1743
A set of analyses that are preserved following a run of a transformation pass.
Definition: PassManager.h:152
static PreservedAnalyses all()
Construct a special preserved set that preserves all passes.
Definition: PassManager.h:158
void preserveSet()
Mark an analysis set as preserved.
Definition: PassManager.h:188
void preserve()
Mark an analysis as preserved.
Definition: PassManager.h:173
static SelectInst * Create(Value *C, Value *S1, Value *S2, const Twine &NameStr="", Instruction *InsertBefore=nullptr, Instruction *MDFrom=nullptr)
A vector that has set insertion semantics.
Definition: SetVector.h:57
This class consists of common code factored out of the SmallVector class to reduce code duplication b...
Definition: SmallVector.h:577
void push_back(const T &Elt)
Definition: SmallVector.h:416
This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.
Definition: SmallVector.h:1200
static unsigned getPointerOperandIndex()
Definition: Instructions.h:395
Analysis pass providing the TargetTransformInfo.
Wrapper pass for TargetTransformInfo.
This pass provides access to the codegen interfaces that are needed for IR-level transformations.
unsigned getAssumedAddrSpace(const Value *V) const
bool isNoopAddrSpaceCast(unsigned FromAS, unsigned ToAS) const
std::pair< const Value *, unsigned > getPredicatedAddrSpace(const Value *V) const
bool collectFlatAddressOperands(SmallVectorImpl< int > &OpIndexes, Intrinsic::ID IID) const
Return any intrinsic address operand indexes which may be rewritten if they use a flat address space ...
Value * rewriteIntrinsicWithAddressSpace(IntrinsicInst *II, Value *OldV, Value *NewV) const
Rewrite intrinsic call II such that OldV will be replaced with NewV, which has a different address sp...
unsigned getFlatAddressSpace() const
Returns the address space ID for a target's 'flat' address space.
bool hasVolatileVariant(Instruction *I, unsigned AddrSpace) const
Return true if the given instruction (assumed to be a memory access instruction) has a volatile varia...
The instances of the Type class are immutable: once they are created, they are never changed.
Definition: Type.h:45
unsigned getPointerAddressSpace() const
Get the address space of this pointer or pointer vector type.
A Use represents the edge between a Value definition and its users.
Definition: Use.h:43
User * getUser() const
Returns the User that contains this Use.
Definition: Use.h:72
Value * get() const
Definition: Use.h:66
void setOperand(unsigned i, Value *Val)
Definition: User.h:174
Value * getOperand(unsigned i) const
Definition: User.h:169
ValueT lookup(const KeyT &Val) const
lookup - Return the entry for the specified key, or a default constructed value if no such entry exis...
Definition: ValueMap.h:164
bool empty() const
Definition: ValueMap.h:139
LLVM Value Representation.
Definition: Value.h:74
Type * getType() const
All values are typed, get the type of this value.
Definition: Value.h:255
void replaceAllUsesWith(Value *V)
Change all uses of this to point to a new Value.
Definition: Value.cpp:535
const Value * stripInBoundsOffsets(function_ref< void(const Value *)> Func=[](const Value *) {}) const
Strip off pointer casts and inbounds GEPs.
Definition: Value.cpp:780
use_iterator_impl< Use > use_iterator
Definition: Value.h:353
Value handle that is nullable, but tries to track the Value.
Definition: ValueHandle.h:204
std::pair< iterator, bool > insert(const ValueT &V)
Definition: DenseSet.h:206
#define llvm_unreachable(msg)
Marks that the current location is not supposed to be reachable.
unsigned ID
LLVM IR allows to use arbitrary numbers as calling convention identifiers.
Definition: CallingConv.h:24
@ C
The default llvm calling convention, compatible with C.
Definition: CallingConv.h:34
Function * getDeclaration(Module *M, ID id, ArrayRef< Type * > Tys=std::nullopt)
Create or insert an LLVM Function declaration for an intrinsic, and return it.
Definition: Function.cpp:1422
@ ReallyHidden
Definition: CommandLine.h:139
initializer< Ty > init(const Ty &Val)
Definition: CommandLine.h:445
PointerTypeMap run(const Module &M)
Compute the PointerTypeMap for the module M.
This is an optimization pass for GlobalISel generic memory operations.
Definition: AddressRanges.h:18
bool RecursivelyDeleteTriviallyDeadInstructions(Value *V, const TargetLibraryInfo *TLI=nullptr, MemorySSAUpdater *MSSAU=nullptr, std::function< void(Value *)> AboutToDeleteCallback=std::function< void(Value *)>())
If the specified value is a trivially dead instruction, delete it.
Definition: Local.cpp:529
AddressSpace
Definition: NVPTXBaseInfo.h:21
void initializeInferAddressSpacesPass(PassRegistry &)
raw_ostream & dbgs()
dbgs() - This returns a reference to a raw_ostream for debugging messages.
Definition: Debug.cpp:163
FunctionPass * createInferAddressSpacesPass(unsigned AddressSpace=~0u)
bool isValidAssumeForContext(const Instruction *I, const Instruction *CxtI, const DominatorTree *DT=nullptr)
Return true if it is valid to use the assumptions provided by an assume intrinsic,...
PreservedAnalyses run(Function &F, FunctionAnalysisManager &AM)