Mercurial > hg > CbC > CbC_llvm

comparison lib/Target/NVPTX/NVPTXFavorNonGenericAddrSpaces.cpp @ 77:54457678186b LLVM3.6

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
LLVM 3.6
author Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp>
date Mon, 08 Sep 2014 22:06:00 +0900
parents
children afa8332a0e37
comparison
equal deleted inserted replaced
34:e874dbf0ad9d 77:54457678186b
1 //===-- NVPTXFavorNonGenericAddrSpace.cpp - ---------------------*- C++ -*-===//
2 //
3 // The LLVM Compiler Infrastructure
4 //
5 // This file is distributed under the University of Illinois Open Source
6 // License. See LICENSE.TXT for details.
7 //
8 //===----------------------------------------------------------------------===//
9 //
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 seeking addrspacecasts, this optimization also traces into
14 // the base pointer of a GEP.
15 //
16 // For instance, the code below loads a float from an array allocated in
17 // addrspace(3).
18 //
19 // %0 = addrspacecast [10 x float] addrspace(3)* @a to [10 x float]*
20 // %1 = gep [10 x float]* %0, i64 0, i64 %i
21 // %2 = load float* %1 ; emits ld.f32
22 //
23 // First, function hoistAddrSpaceCastFromGEP reorders the addrspacecast
24 // and the GEP to expose more optimization opportunities to function
25 // optimizeMemoryInst. The intermediate code looks like:
26 //
27 // %0 = gep [10 x float] addrspace(3)* @a, i64 0, i64 %i
28 // %1 = addrspacecast float addrspace(3)* %0 to float*
29 // %2 = load float* %1 ; still emits ld.f32, but will be optimized shortly
30 //
31 // Then, function optimizeMemoryInstruction detects a load from addrspacecast'ed
32 // generic pointers, and folds the load and the addrspacecast into a load from
33 // the original address space. The final code looks like:
34 //
35 // %0 = gep [10 x float] addrspace(3)* @a, i64 0, i64 %i
36 // %2 = load float addrspace(3)* %0 ; emits ld.shared.f32
37 //
38 // This pass may remove an addrspacecast in a different BB. Therefore, we
39 // implement it as a FunctionPass.
40 //
41 //===----------------------------------------------------------------------===//
42
43 #include "NVPTX.h"
44 #include "llvm/IR/Function.h"
45 #include "llvm/IR/Instructions.h"
46 #include "llvm/IR/Operator.h"
47 #include "llvm/Support/CommandLine.h"
48
49 using namespace llvm;
50
51 // An option to disable this optimization. Enable it by default.
52 static cl::opt<bool> DisableFavorNonGeneric(
53 "disable-nvptx-favor-non-generic",
54 cl::init(false),
55 cl::desc("Do not convert generic address space usage "
56 "to non-generic address space usage"),
57 cl::Hidden);
58
59 namespace {
60 /// \brief NVPTXFavorNonGenericAddrSpaces
61 class NVPTXFavorNonGenericAddrSpaces : public FunctionPass {
62 public:
63 static char ID;
64 NVPTXFavorNonGenericAddrSpaces() : FunctionPass(ID) {}
65
66 bool runOnFunction(Function &F) override;
67
68 /// Optimizes load/store instructions. Idx is the index of the pointer operand
69 /// (0 for load, and 1 for store). Returns true if it changes anything.
70 bool optimizeMemoryInstruction(Instruction *I, unsigned Idx);
71 /// Transforms "gep (addrspacecast X), indices" into "addrspacecast (gep X,
72 /// indices)". This reordering exposes to optimizeMemoryInstruction more
73 /// optimization opportunities on loads and stores. Returns true if it changes
74 /// the program.
75 bool hoistAddrSpaceCastFromGEP(GEPOperator *GEP);
76 };
77 }
78
79 char NVPTXFavorNonGenericAddrSpaces::ID = 0;
80
81 namespace llvm {
82 void initializeNVPTXFavorNonGenericAddrSpacesPass(PassRegistry &);
83 }
84 INITIALIZE_PASS(NVPTXFavorNonGenericAddrSpaces, "nvptx-favor-non-generic",
85 "Remove unnecessary non-generic-to-generic addrspacecasts",
86 false, false)
87
88 // Decides whether removing Cast is valid and beneficial. Cast can be an
89 // instruction or a constant expression.
90 static bool IsEliminableAddrSpaceCast(Operator *Cast) {
91 // Returns false if not even an addrspacecast.
92 if (Cast->getOpcode() != Instruction::AddrSpaceCast)
93 return false;
94
95 Value *Src = Cast->getOperand(0);
96 PointerType *SrcTy = cast<PointerType>(Src->getType());
97 PointerType *DestTy = cast<PointerType>(Cast->getType());
98 // TODO: For now, we only handle the case where the addrspacecast only changes
99 // the address space but not the type. If the type also changes, we could
100 // still get rid of the addrspacecast by adding an extra bitcast, but we
101 // rarely see such scenarios.
102 if (SrcTy->getElementType() != DestTy->getElementType())
103 return false;
104
105 // Checks whether the addrspacecast is from a non-generic address space to the
106 // generic address space.
107 return (SrcTy->getAddressSpace() != AddressSpace::ADDRESS_SPACE_GENERIC &&
108 DestTy->getAddressSpace() == AddressSpace::ADDRESS_SPACE_GENERIC);
109 }
110
111 bool NVPTXFavorNonGenericAddrSpaces::hoistAddrSpaceCastFromGEP(
112 GEPOperator *GEP) {
113 Operator *Cast = dyn_cast<Operator>(GEP->getPointerOperand());
114 if (!Cast)
115 return false;
116
117 if (!IsEliminableAddrSpaceCast(Cast))
118 return false;
119
120 SmallVector<Value *, 8> Indices(GEP->idx_begin(), GEP->idx_end());
121 if (Instruction *GEPI = dyn_cast<Instruction>(GEP)) {
122 // %1 = gep (addrspacecast X), indices
123 // =>
124 // %0 = gep X, indices
125 // %1 = addrspacecast %0
126 GetElementPtrInst *NewGEPI = GetElementPtrInst::Create(Cast->getOperand(0),
127 Indices,
128 GEP->getName(),
129 GEPI);
130 NewGEPI->setIsInBounds(GEP->isInBounds());
131 GEP->replaceAllUsesWith(
132 new AddrSpaceCastInst(NewGEPI, GEP->getType(), "", GEPI));
133 } else {
134 // GEP is a constant expression.
135 Constant *NewGEPCE = ConstantExpr::getGetElementPtr(
136 cast<Constant>(Cast->getOperand(0)),
137 Indices,
138 GEP->isInBounds());
139 GEP->replaceAllUsesWith(
140 ConstantExpr::getAddrSpaceCast(NewGEPCE, GEP->getType()));
141 }
142
143 return true;
144 }
145
146 bool NVPTXFavorNonGenericAddrSpaces::optimizeMemoryInstruction(Instruction *MI,
147 unsigned Idx) {
148 // If the pointer operand is a GEP, hoist the addrspacecast if any from the
149 // GEP to expose more optimization opportunites.
150 if (GEPOperator *GEP = dyn_cast<GEPOperator>(MI->getOperand(Idx))) {
151 hoistAddrSpaceCastFromGEP(GEP);
152 }
153
154 // load/store (addrspacecast X) => load/store X if shortcutting the
155 // addrspacecast is valid and can improve performance.
156 //
157 // e.g.,
158 // %1 = addrspacecast float addrspace(3)* %0 to float*
159 // %2 = load float* %1
160 // ->
161 // %2 = load float addrspace(3)* %0
162 //
163 // Note: the addrspacecast can also be a constant expression.
164 if (Operator *Cast = dyn_cast<Operator>(MI->getOperand(Idx))) {
165 if (IsEliminableAddrSpaceCast(Cast)) {
166 MI->setOperand(Idx, Cast->getOperand(0));
167 return true;
168 }
169 }
170
171 return false;
172 }
173
174 bool NVPTXFavorNonGenericAddrSpaces::runOnFunction(Function &F) {
175 if (DisableFavorNonGeneric)
176 return false;
177
178 bool Changed = false;
179 for (Function::iterator B = F.begin(), BE = F.end(); B != BE; ++B) {
180 for (BasicBlock::iterator I = B->begin(), IE = B->end(); I != IE; ++I) {
181 if (isa<LoadInst>(I)) {
182 // V = load P
183 Changed |= optimizeMemoryInstruction(I, 0);
184 } else if (isa<StoreInst>(I)) {
185 // store V, P
186 Changed |= optimizeMemoryInstruction(I, 1);
187 }
188 }
189 }
190 return Changed;
191 }
192
193 FunctionPass *llvm::createNVPTXFavorNonGenericAddrSpacesPass() {
194 return new NVPTXFavorNonGenericAddrSpaces();
195 }