1 //===-- NVPTXFavorNonGenericAddrSpace.cpp - ---------------------*- C++ -*-===//
3 // The LLVM Compiler Infrastructure
5 // This file is distributed under the University of Illinois Open Source
6 // License. See LICENSE.TXT for details.
8 //===----------------------------------------------------------------------===//
10 // When a load/store accesses the generic address space, checks whether the
11 // address is casted from a non-generic address space. If so, remove this
12 // addrspacecast because accessing non-generic address spaces is typically
13 // faster. Besides removing addrspacecasts directly used by loads/stores, this
14 // optimization also recursively traces into a GEP's pointer operand and a
15 // bitcast's source to find more eliminable addrspacecasts.
17 // For instance, the code below loads a float from an array allocated in
20 // %0 = addrspacecast [10 x float] addrspace(3)* @a to [10 x float]*
21 // %1 = gep [10 x float]* %0, i64 0, i64 %i
22 // %2 = bitcast float* %1 to i32*
23 // %3 = load i32* %2 ; emits ld.u32
25 // First, function hoistAddrSpaceCastFrom reorders the addrspacecast, the GEP,
26 // and the bitcast to expose more optimization opportunities to function
27 // optimizeMemoryInst. The intermediate code looks like:
29 // %0 = gep [10 x float] addrspace(3)* @a, i64 0, i64 %i
30 // %1 = bitcast float addrspace(3)* %0 to i32 addrspace(3)*
31 // %2 = addrspacecast i32 addrspace(3)* %1 to i32*
32 // %3 = load i32* %2 ; still emits ld.u32, but will be optimized shortly
34 // Then, function optimizeMemoryInstruction detects a load from addrspacecast'ed
35 // generic pointers, and folds the load and the addrspacecast into a load from
36 // the original address space. The final code looks like:
38 // %0 = gep [10 x float] addrspace(3)* @a, i64 0, i64 %i
39 // %1 = bitcast float addrspace(3)* %0 to i32 addrspace(3)*
40 // %3 = load i32 addrspace(3)* %1 ; emits ld.shared.f32
42 // This pass may remove an addrspacecast in a different BB. Therefore, we
43 // implement it as a FunctionPass.
46 // The current implementation doesn't handle PHINodes. Eliminating
47 // addrspacecasts used by PHINodes is trickier because PHINodes can introduce
48 // loops in data flow. For example,
50 // %generic.input = addrspacecast float addrspace(3)* %input to float*
52 // %y = phi [ %generic.input, %y2 ]
53 // %y2 = getelementptr %y, 1
55 // br ..., label %loop, ...
57 // Marking %y2 shared depends on marking %y shared, but %y also data-flow
58 // depends on %y2. We probably need an iterative fix-point algorithm on handle
61 //===----------------------------------------------------------------------===//
64 #include "llvm/IR/Function.h"
65 #include "llvm/IR/Instructions.h"
66 #include "llvm/IR/Operator.h"
67 #include "llvm/Support/CommandLine.h"
71 // An option to disable this optimization. Enable it by default.
72 static cl::opt<bool> DisableFavorNonGeneric(
73 "disable-nvptx-favor-non-generic",
75 cl::desc("Do not convert generic address space usage "
76 "to non-generic address space usage"),
80 /// \brief NVPTXFavorNonGenericAddrSpaces
81 class NVPTXFavorNonGenericAddrSpaces : public FunctionPass {
84 NVPTXFavorNonGenericAddrSpaces() : FunctionPass(ID) {}
85 bool runOnFunction(Function &F) override;
88 /// Optimizes load/store instructions. Idx is the index of the pointer operand
89 /// (0 for load, and 1 for store). Returns true if it changes anything.
90 bool optimizeMemoryInstruction(Instruction *I, unsigned Idx);
91 /// Recursively traces into a GEP's pointer operand or a bitcast's source to
92 /// find an eliminable addrspacecast, and hoists that addrspacecast to the
93 /// outermost level. For example, this function transforms
94 /// bitcast(gep(gep(addrspacecast(X))))
96 /// addrspacecast(bitcast(gep(gep(X)))).
98 /// This reordering exposes to optimizeMemoryInstruction more
99 /// optimization opportunities on loads and stores.
101 /// Returns true if this function succesfully hoists an eliminable
102 /// addrspacecast or V is already such an addrspacecast.
103 /// Transforms "gep (addrspacecast X), indices" into "addrspacecast (gep X,
105 bool hoistAddrSpaceCastFrom(Value *V, int Depth = 0);
106 /// Helper function for GEPs.
107 bool hoistAddrSpaceCastFromGEP(GEPOperator *GEP, int Depth);
108 /// Helper function for bitcasts.
109 bool hoistAddrSpaceCastFromBitCast(BitCastOperator *BC, int Depth);
113 char NVPTXFavorNonGenericAddrSpaces::ID = 0;
116 void initializeNVPTXFavorNonGenericAddrSpacesPass(PassRegistry &);
118 INITIALIZE_PASS(NVPTXFavorNonGenericAddrSpaces, "nvptx-favor-non-generic",
119 "Remove unnecessary non-generic-to-generic addrspacecasts",
122 // Decides whether V is an addrspacecast and shortcutting V in load/store is
123 // valid and beneficial.
124 static bool isEliminableAddrSpaceCast(Value *V) {
125 // Returns false if V is not even an addrspacecast.
126 Operator *Cast = dyn_cast<Operator>(V);
127 if (Cast == nullptr || Cast->getOpcode() != Instruction::AddrSpaceCast)
130 Value *Src = Cast->getOperand(0);
131 PointerType *SrcTy = cast<PointerType>(Src->getType());
132 PointerType *DestTy = cast<PointerType>(Cast->getType());
133 // TODO: For now, we only handle the case where the addrspacecast only changes
134 // the address space but not the type. If the type also changes, we could
135 // still get rid of the addrspacecast by adding an extra bitcast, but we
136 // rarely see such scenarios.
137 if (SrcTy->getElementType() != DestTy->getElementType())
140 // Checks whether the addrspacecast is from a non-generic address space to the
141 // generic address space.
142 return (SrcTy->getAddressSpace() != AddressSpace::ADDRESS_SPACE_GENERIC &&
143 DestTy->getAddressSpace() == AddressSpace::ADDRESS_SPACE_GENERIC);
146 bool NVPTXFavorNonGenericAddrSpaces::hoistAddrSpaceCastFromGEP(GEPOperator *GEP,
148 if (!hoistAddrSpaceCastFrom(GEP->getPointerOperand(), Depth + 1))
151 // That hoistAddrSpaceCastFrom succeeds implies GEP's pointer operand is now
152 // an eliminable addrspacecast.
153 assert(isEliminableAddrSpaceCast(GEP->getPointerOperand()));
154 Operator *Cast = cast<Operator>(GEP->getPointerOperand());
156 SmallVector<Value *, 8> Indices(GEP->idx_begin(), GEP->idx_end());
157 if (Instruction *GEPI = dyn_cast<Instruction>(GEP)) {
158 // GEP = gep (addrspacecast X), indices
160 // NewGEP = gep X, indices
161 // NewASC = addrspacecast NewGEP
162 GetElementPtrInst *NewGEP = GetElementPtrInst::Create(
163 GEP->getSourceElementType(), Cast->getOperand(0), Indices,
165 NewGEP->setIsInBounds(GEP->isInBounds());
166 Value *NewASC = new AddrSpaceCastInst(NewGEP, GEP->getType(), "", GEPI);
167 NewASC->takeName(GEP);
168 GEP->replaceAllUsesWith(NewASC);
170 // GEP is a constant expression.
171 Constant *NewGEP = ConstantExpr::getGetElementPtr(
172 GEP->getSourceElementType(), cast<Constant>(Cast->getOperand(0)),
173 Indices, GEP->isInBounds());
174 GEP->replaceAllUsesWith(
175 ConstantExpr::getAddrSpaceCast(NewGEP, GEP->getType()));
181 bool NVPTXFavorNonGenericAddrSpaces::hoistAddrSpaceCastFromBitCast(
182 BitCastOperator *BC, int Depth) {
183 if (!hoistAddrSpaceCastFrom(BC->getOperand(0), Depth + 1))
186 // That hoistAddrSpaceCastFrom succeeds implies BC's source operand is now
187 // an eliminable addrspacecast.
188 assert(isEliminableAddrSpaceCast(BC->getOperand(0)));
189 Operator *Cast = cast<Operator>(BC->getOperand(0));
191 // Cast = addrspacecast Src
194 // Cast' = bitcast Src
195 // BC' = addrspacecast Cast'
196 Value *Src = Cast->getOperand(0);
197 Type *TypeOfNewCast =
198 PointerType::get(BC->getType()->getPointerElementType(),
199 Src->getType()->getPointerAddressSpace());
200 if (BitCastInst *BCI = dyn_cast<BitCastInst>(BC)) {
201 Value *NewCast = new BitCastInst(Src, TypeOfNewCast, "", BCI);
202 Value *NewBC = new AddrSpaceCastInst(NewCast, BC->getType(), "", BCI);
204 BC->replaceAllUsesWith(NewBC);
206 // BC is a constant expression.
208 ConstantExpr::getBitCast(cast<Constant>(Src), TypeOfNewCast);
209 Constant *NewBC = ConstantExpr::getAddrSpaceCast(NewCast, BC->getType());
210 BC->replaceAllUsesWith(NewBC);
215 bool NVPTXFavorNonGenericAddrSpaces::hoistAddrSpaceCastFrom(Value *V,
217 // Returns true if V is already an eliminable addrspacecast.
218 if (isEliminableAddrSpaceCast(V))
221 // Limit the depth to prevent this recursive function from running too long.
222 const int MaxDepth = 20;
223 if (Depth >= MaxDepth)
226 // If V is a GEP or bitcast, hoist the addrspacecast if any from its pointer
227 // operand. This enables optimizeMemoryInstruction to shortcut addrspacecasts
228 // that are not directly used by the load/store.
229 if (GEPOperator *GEP = dyn_cast<GEPOperator>(V))
230 return hoistAddrSpaceCastFromGEP(GEP, Depth);
232 if (BitCastOperator *BC = dyn_cast<BitCastOperator>(V))
233 return hoistAddrSpaceCastFromBitCast(BC, Depth);
238 bool NVPTXFavorNonGenericAddrSpaces::optimizeMemoryInstruction(Instruction *MI,
240 if (hoistAddrSpaceCastFrom(MI->getOperand(Idx))) {
241 // load/store (addrspacecast X) => load/store X if shortcutting the
242 // addrspacecast is valid and can improve performance.
245 // %1 = addrspacecast float addrspace(3)* %0 to float*
246 // %2 = load float* %1
248 // %2 = load float addrspace(3)* %0
250 // Note: the addrspacecast can also be a constant expression.
251 assert(isEliminableAddrSpaceCast(MI->getOperand(Idx)));
252 Operator *ASC = dyn_cast<Operator>(MI->getOperand(Idx));
253 MI->setOperand(Idx, ASC->getOperand(0));
259 bool NVPTXFavorNonGenericAddrSpaces::runOnFunction(Function &F) {
260 if (DisableFavorNonGeneric)
263 bool Changed = false;
264 for (Function::iterator B = F.begin(), BE = F.end(); B != BE; ++B) {
265 for (BasicBlock::iterator I = B->begin(), IE = B->end(); I != IE; ++I) {
266 if (isa<LoadInst>(I)) {
268 Changed |= optimizeMemoryInstruction(I, 0);
269 } else if (isa<StoreInst>(I)) {
271 Changed |= optimizeMemoryInstruction(I, 1);
278 FunctionPass *llvm::createNVPTXFavorNonGenericAddrSpacesPass() {
279 return new NVPTXFavorNonGenericAddrSpaces();