CbC/CbC_llvm: lib/Target/NVPTX/NVPTXFavorNonGenericAddrSpaces.cpp annotate

author	Kaito Tokumori <e105711@ie.u-ryukyu.ac.jp>
date	Mon, 08 Sep 2014 22:06:00 +0900
parents
children	afa8332a0e37

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

77

54457678186b LLVM 3.6