NVPTXFavorNonGenericAddrSpaces.cpp revision 4c5e43da7792f75567b693105cc53e3f1992ad98
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 49using namespace llvm; 50 51// An option to disable this optimization. Enable it by default. 52static 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 59namespace { 60/// \brief NVPTXFavorNonGenericAddrSpaces 61class NVPTXFavorNonGenericAddrSpaces : public FunctionPass { 62public: 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 79char NVPTXFavorNonGenericAddrSpaces::ID = 0; 80 81namespace llvm { 82void initializeNVPTXFavorNonGenericAddrSpacesPass(PassRegistry &); 83} 84INITIALIZE_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. 90static 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 111bool 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( 127 GEP->getSourceElementType(), Cast->getOperand(0), Indices, 128 GEP->getName(), GEPI); 129 NewGEPI->setIsInBounds(GEP->isInBounds()); 130 GEP->replaceAllUsesWith( 131 new AddrSpaceCastInst(NewGEPI, GEP->getType(), "", GEPI)); 132 } else { 133 // GEP is a constant expression. 134 Constant *NewGEPCE = ConstantExpr::getGetElementPtr( 135 cast<Constant>(Cast->getOperand(0)), 136 Indices, 137 GEP->isInBounds()); 138 GEP->replaceAllUsesWith( 139 ConstantExpr::getAddrSpaceCast(NewGEPCE, GEP->getType())); 140 } 141 142 return true; 143} 144 145bool NVPTXFavorNonGenericAddrSpaces::optimizeMemoryInstruction(Instruction *MI, 146 unsigned Idx) { 147 // If the pointer operand is a GEP, hoist the addrspacecast if any from the 148 // GEP to expose more optimization opportunites. 149 if (GEPOperator *GEP = dyn_cast<GEPOperator>(MI->getOperand(Idx))) { 150 hoistAddrSpaceCastFromGEP(GEP); 151 } 152 153 // load/store (addrspacecast X) => load/store X if shortcutting the 154 // addrspacecast is valid and can improve performance. 155 // 156 // e.g., 157 // %1 = addrspacecast float addrspace(3)* %0 to float* 158 // %2 = load float* %1 159 // -> 160 // %2 = load float addrspace(3)* %0 161 // 162 // Note: the addrspacecast can also be a constant expression. 163 if (Operator *Cast = dyn_cast<Operator>(MI->getOperand(Idx))) { 164 if (IsEliminableAddrSpaceCast(Cast)) { 165 MI->setOperand(Idx, Cast->getOperand(0)); 166 return true; 167 } 168 } 169 170 return false; 171} 172 173bool NVPTXFavorNonGenericAddrSpaces::runOnFunction(Function &F) { 174 if (DisableFavorNonGeneric) 175 return false; 176 177 bool Changed = false; 178 for (Function::iterator B = F.begin(), BE = F.end(); B != BE; ++B) { 179 for (BasicBlock::iterator I = B->begin(), IE = B->end(); I != IE; ++I) { 180 if (isa<LoadInst>(I)) { 181 // V = load P 182 Changed |= optimizeMemoryInstruction(I, 0); 183 } else if (isa<StoreInst>(I)) { 184 // store V, P 185 Changed |= optimizeMemoryInstruction(I, 1); 186 } 187 } 188 } 189 return Changed; 190} 191 192FunctionPass *llvm::createNVPTXFavorNonGenericAddrSpacesPass() { 193 return new NVPTXFavorNonGenericAddrSpaces(); 194} 195