lib/Target/NVPTX/NVPTXFavorNonGenericAddrSpaces.cpp

   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(
 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         GEP->getSourceElementType(), cast<Constant>(Cast->getOperand(0)),
 136         Indices, GEP->isInBounds());
 137     GEP->replaceAllUsesWith(
 138         ConstantExpr::getAddrSpaceCast(NewGEPCE, GEP->getType()));
 139   }
 140
 141   return true;
 142 }
 143
 144 bool NVPTXFavorNonGenericAddrSpaces::optimizeMemoryInstruction(Instruction *MI,
 145                                                                unsigned Idx) {
 146   // If the pointer operand is a GEP, hoist the addrspacecast if any from the
 147   // GEP to expose more optimization opportunites.
 148   if (GEPOperator *GEP = dyn_cast<GEPOperator>(MI->getOperand(Idx))) {
 149     hoistAddrSpaceCastFromGEP(GEP);
 150   }
 151
 152   // load/store (addrspacecast X) => load/store X if shortcutting the
 153   // addrspacecast is valid and can improve performance.
 154   //
 155   // e.g.,
 156   // %1 = addrspacecast float addrspace(3)* %0 to float*
 157   // %2 = load float* %1
 158   // ->
 159   // %2 = load float addrspace(3)* %0
 160   //
 161   // Note: the addrspacecast can also be a constant expression.
 162   if (Operator *Cast = dyn_cast<Operator>(MI->getOperand(Idx))) {
 163     if (IsEliminableAddrSpaceCast(Cast)) {
 164       MI->setOperand(Idx, Cast->getOperand(0));
 165       return true;
 166     }
 167   }
 168
 169   return false;
 170 }
 171
 172 bool NVPTXFavorNonGenericAddrSpaces::runOnFunction(Function &F) {
 173   if (DisableFavorNonGeneric)
 174     return false;
 175
 176   bool Changed = false;
 177   for (Function::iterator B = F.begin(), BE = F.end(); B != BE; ++B) {
 178     for (BasicBlock::iterator I = B->begin(), IE = B->end(); I != IE; ++I) {
 179       if (isa<LoadInst>(I)) {
 180         // V = load P
 181         Changed |= optimizeMemoryInstruction(I, 0);
 182       } else if (isa<StoreInst>(I)) {
 183         // store V, P
 184         Changed |= optimizeMemoryInstruction(I, 1);
 185       }
 186     }
 187   }
 188   return Changed;
 189 }
 190
 191 FunctionPass *llvm::createNVPTXFavorNonGenericAddrSpacesPass() {
 192   return new NVPTXFavorNonGenericAddrSpaces();
 193 }