2014-04-04 05:18:25 +08:00
|
|
|
//===-- NVPTXFavorNonGenericAddrSpace.cpp - ---------------------*- C++ -*-===//
|
|
|
|
//
|
|
|
|
// The LLVM Compiler Infrastructure
|
|
|
|
//
|
|
|
|
// This file is distributed under the University of Illinois Open Source
|
|
|
|
// License. See LICENSE.TXT for details.
|
|
|
|
//
|
|
|
|
//===----------------------------------------------------------------------===//
|
|
|
|
//
|
2016-03-21 04:59:20 +08:00
|
|
|
// FIXME: This pass is deprecated in favor of NVPTXInferAddressSpaces, which
|
|
|
|
// uses a new algorithm that handles pointer induction variables.
|
|
|
|
//
|
2014-04-04 05:18:25 +08:00
|
|
|
// When a load/store accesses the generic address space, checks whether the
|
|
|
|
// address is casted from a non-generic address space. If so, remove this
|
|
|
|
// addrspacecast because accessing non-generic address spaces is typically
|
2015-05-30 01:00:27 +08:00
|
|
|
// faster. Besides removing addrspacecasts directly used by loads/stores, this
|
|
|
|
// optimization also recursively traces into a GEP's pointer operand and a
|
|
|
|
// bitcast's source to find more eliminable addrspacecasts.
|
2014-04-04 05:18:25 +08:00
|
|
|
//
|
|
|
|
// For instance, the code below loads a float from an array allocated in
|
|
|
|
// addrspace(3).
|
|
|
|
//
|
2015-05-30 01:00:27 +08:00
|
|
|
// %0 = addrspacecast [10 x float] addrspace(3)* @a to [10 x float]*
|
|
|
|
// %1 = gep [10 x float]* %0, i64 0, i64 %i
|
|
|
|
// %2 = bitcast float* %1 to i32*
|
|
|
|
// %3 = load i32* %2 ; emits ld.u32
|
2014-04-04 05:18:25 +08:00
|
|
|
//
|
2015-05-30 01:00:27 +08:00
|
|
|
// First, function hoistAddrSpaceCastFrom reorders the addrspacecast, the GEP,
|
|
|
|
// and the bitcast to expose more optimization opportunities to function
|
2014-04-04 05:18:25 +08:00
|
|
|
// optimizeMemoryInst. The intermediate code looks like:
|
|
|
|
//
|
2015-05-30 01:00:27 +08:00
|
|
|
// %0 = gep [10 x float] addrspace(3)* @a, i64 0, i64 %i
|
|
|
|
// %1 = bitcast float addrspace(3)* %0 to i32 addrspace(3)*
|
|
|
|
// %2 = addrspacecast i32 addrspace(3)* %1 to i32*
|
|
|
|
// %3 = load i32* %2 ; still emits ld.u32, but will be optimized shortly
|
2014-04-04 05:18:25 +08:00
|
|
|
//
|
|
|
|
// Then, function optimizeMemoryInstruction detects a load from addrspacecast'ed
|
|
|
|
// generic pointers, and folds the load and the addrspacecast into a load from
|
|
|
|
// the original address space. The final code looks like:
|
|
|
|
//
|
2015-05-30 01:00:27 +08:00
|
|
|
// %0 = gep [10 x float] addrspace(3)* @a, i64 0, i64 %i
|
|
|
|
// %1 = bitcast float addrspace(3)* %0 to i32 addrspace(3)*
|
|
|
|
// %3 = load i32 addrspace(3)* %1 ; emits ld.shared.f32
|
2014-04-04 05:18:25 +08:00
|
|
|
//
|
|
|
|
// This pass may remove an addrspacecast in a different BB. Therefore, we
|
|
|
|
// implement it as a FunctionPass.
|
|
|
|
//
|
2015-05-30 01:00:27 +08:00
|
|
|
// TODO:
|
|
|
|
// The current implementation doesn't handle PHINodes. Eliminating
|
|
|
|
// addrspacecasts used by PHINodes is trickier because PHINodes can introduce
|
|
|
|
// loops in data flow. For example,
|
|
|
|
//
|
|
|
|
// %generic.input = addrspacecast float addrspace(3)* %input to float*
|
|
|
|
// loop:
|
|
|
|
// %y = phi [ %generic.input, %y2 ]
|
|
|
|
// %y2 = getelementptr %y, 1
|
|
|
|
// %v = load %y2
|
|
|
|
// br ..., label %loop, ...
|
|
|
|
//
|
|
|
|
// Marking %y2 shared depends on marking %y shared, but %y also data-flow
|
|
|
|
// depends on %y2. We probably need an iterative fix-point algorithm on handle
|
|
|
|
// this case.
|
|
|
|
//
|
2014-04-04 05:18:25 +08:00
|
|
|
//===----------------------------------------------------------------------===//
|
|
|
|
|
|
|
|
#include "NVPTX.h"
|
|
|
|
#include "llvm/IR/Function.h"
|
|
|
|
#include "llvm/IR/Instructions.h"
|
|
|
|
#include "llvm/IR/Operator.h"
|
|
|
|
#include "llvm/Support/CommandLine.h"
|
|
|
|
|
|
|
|
using namespace llvm;
|
|
|
|
|
|
|
|
// An option to disable this optimization. Enable it by default.
|
|
|
|
static cl::opt<bool> DisableFavorNonGeneric(
|
|
|
|
"disable-nvptx-favor-non-generic",
|
|
|
|
cl::init(false),
|
|
|
|
cl::desc("Do not convert generic address space usage "
|
|
|
|
"to non-generic address space usage"),
|
|
|
|
cl::Hidden);
|
|
|
|
|
|
|
|
namespace {
|
|
|
|
/// \brief NVPTXFavorNonGenericAddrSpaces
|
|
|
|
class NVPTXFavorNonGenericAddrSpaces : public FunctionPass {
|
|
|
|
public:
|
|
|
|
static char ID;
|
|
|
|
NVPTXFavorNonGenericAddrSpaces() : FunctionPass(ID) {}
|
2014-05-23 14:30:12 +08:00
|
|
|
bool runOnFunction(Function &F) override;
|
2014-04-04 05:18:25 +08:00
|
|
|
|
2015-05-30 01:00:27 +08:00
|
|
|
private:
|
2014-04-04 05:18:25 +08:00
|
|
|
/// Optimizes load/store instructions. Idx is the index of the pointer operand
|
|
|
|
/// (0 for load, and 1 for store). Returns true if it changes anything.
|
|
|
|
bool optimizeMemoryInstruction(Instruction *I, unsigned Idx);
|
2015-05-30 01:00:27 +08:00
|
|
|
/// Recursively traces into a GEP's pointer operand or a bitcast's source to
|
|
|
|
/// find an eliminable addrspacecast, and hoists that addrspacecast to the
|
|
|
|
/// outermost level. For example, this function transforms
|
|
|
|
/// bitcast(gep(gep(addrspacecast(X))))
|
|
|
|
/// to
|
|
|
|
/// addrspacecast(bitcast(gep(gep(X)))).
|
|
|
|
///
|
|
|
|
/// This reordering exposes to optimizeMemoryInstruction more
|
|
|
|
/// optimization opportunities on loads and stores.
|
|
|
|
///
|
2015-08-09 02:27:36 +08:00
|
|
|
/// If this function successfully hoists an eliminable addrspacecast or V is
|
2015-06-10 05:50:32 +08:00
|
|
|
/// already such an addrspacecast, it returns the transformed value (which is
|
|
|
|
/// guaranteed to be an addrspacecast); otherwise, it returns nullptr.
|
|
|
|
Value *hoistAddrSpaceCastFrom(Value *V, int Depth = 0);
|
2015-05-30 01:00:27 +08:00
|
|
|
/// Helper function for GEPs.
|
2015-06-10 05:50:32 +08:00
|
|
|
Value *hoistAddrSpaceCastFromGEP(GEPOperator *GEP, int Depth);
|
2015-05-30 01:00:27 +08:00
|
|
|
/// Helper function for bitcasts.
|
2015-06-10 05:50:32 +08:00
|
|
|
Value *hoistAddrSpaceCastFromBitCast(BitCastOperator *BC, int Depth);
|
2014-04-04 05:18:25 +08:00
|
|
|
};
|
2015-06-23 17:49:53 +08:00
|
|
|
}
|
2014-04-04 05:18:25 +08:00
|
|
|
|
|
|
|
char NVPTXFavorNonGenericAddrSpaces::ID = 0;
|
|
|
|
|
|
|
|
namespace llvm {
|
|
|
|
void initializeNVPTXFavorNonGenericAddrSpacesPass(PassRegistry &);
|
|
|
|
}
|
|
|
|
INITIALIZE_PASS(NVPTXFavorNonGenericAddrSpaces, "nvptx-favor-non-generic",
|
|
|
|
"Remove unnecessary non-generic-to-generic addrspacecasts",
|
|
|
|
false, false)
|
|
|
|
|
2015-05-30 01:00:27 +08:00
|
|
|
// Decides whether V is an addrspacecast and shortcutting V in load/store is
|
|
|
|
// valid and beneficial.
|
|
|
|
static bool isEliminableAddrSpaceCast(Value *V) {
|
|
|
|
// Returns false if V is not even an addrspacecast.
|
|
|
|
Operator *Cast = dyn_cast<Operator>(V);
|
|
|
|
if (Cast == nullptr || Cast->getOpcode() != Instruction::AddrSpaceCast)
|
2014-04-04 05:18:25 +08:00
|
|
|
return false;
|
|
|
|
|
|
|
|
Value *Src = Cast->getOperand(0);
|
|
|
|
PointerType *SrcTy = cast<PointerType>(Src->getType());
|
|
|
|
PointerType *DestTy = cast<PointerType>(Cast->getType());
|
|
|
|
// TODO: For now, we only handle the case where the addrspacecast only changes
|
|
|
|
// the address space but not the type. If the type also changes, we could
|
|
|
|
// still get rid of the addrspacecast by adding an extra bitcast, but we
|
|
|
|
// rarely see such scenarios.
|
|
|
|
if (SrcTy->getElementType() != DestTy->getElementType())
|
|
|
|
return false;
|
|
|
|
|
|
|
|
// Checks whether the addrspacecast is from a non-generic address space to the
|
|
|
|
// generic address space.
|
|
|
|
return (SrcTy->getAddressSpace() != AddressSpace::ADDRESS_SPACE_GENERIC &&
|
|
|
|
DestTy->getAddressSpace() == AddressSpace::ADDRESS_SPACE_GENERIC);
|
|
|
|
}
|
|
|
|
|
2015-06-10 05:50:32 +08:00
|
|
|
Value *NVPTXFavorNonGenericAddrSpaces::hoistAddrSpaceCastFromGEP(
|
|
|
|
GEPOperator *GEP, int Depth) {
|
|
|
|
Value *NewOperand =
|
|
|
|
hoistAddrSpaceCastFrom(GEP->getPointerOperand(), Depth + 1);
|
|
|
|
if (NewOperand == nullptr)
|
|
|
|
return nullptr;
|
2014-04-04 05:18:25 +08:00
|
|
|
|
2015-06-10 05:50:32 +08:00
|
|
|
// hoistAddrSpaceCastFrom returns an eliminable addrspacecast or nullptr.
|
|
|
|
assert(isEliminableAddrSpaceCast(NewOperand));
|
|
|
|
Operator *Cast = cast<Operator>(NewOperand);
|
2014-04-04 05:18:25 +08:00
|
|
|
|
|
|
|
SmallVector<Value *, 8> Indices(GEP->idx_begin(), GEP->idx_end());
|
2015-06-10 05:50:32 +08:00
|
|
|
Value *NewASC;
|
2014-04-04 05:18:25 +08:00
|
|
|
if (Instruction *GEPI = dyn_cast<Instruction>(GEP)) {
|
2015-05-30 01:00:27 +08:00
|
|
|
// GEP = gep (addrspacecast X), indices
|
2014-04-04 05:18:25 +08:00
|
|
|
// =>
|
2015-05-30 01:00:27 +08:00
|
|
|
// NewGEP = gep X, indices
|
|
|
|
// NewASC = addrspacecast NewGEP
|
|
|
|
GetElementPtrInst *NewGEP = GetElementPtrInst::Create(
|
2015-03-15 03:53:33 +08:00
|
|
|
GEP->getSourceElementType(), Cast->getOperand(0), Indices,
|
2015-05-30 01:00:27 +08:00
|
|
|
"", GEPI);
|
|
|
|
NewGEP->setIsInBounds(GEP->isInBounds());
|
2016-03-21 04:59:20 +08:00
|
|
|
NewGEP->takeName(GEP);
|
2015-06-10 05:50:32 +08:00
|
|
|
NewASC = new AddrSpaceCastInst(NewGEP, GEP->getType(), "", GEPI);
|
|
|
|
// Without RAUWing GEP, the compiler would visit GEP again and emit
|
|
|
|
// redundant instructions. This is exercised in test @rauw in
|
|
|
|
// access-non-generic.ll.
|
2015-05-30 01:00:27 +08:00
|
|
|
GEP->replaceAllUsesWith(NewASC);
|
2014-04-04 05:18:25 +08:00
|
|
|
} else {
|
|
|
|
// GEP is a constant expression.
|
2015-05-30 01:00:27 +08:00
|
|
|
Constant *NewGEP = ConstantExpr::getGetElementPtr(
|
2015-04-03 02:55:32 +08:00
|
|
|
GEP->getSourceElementType(), cast<Constant>(Cast->getOperand(0)),
|
|
|
|
Indices, GEP->isInBounds());
|
2015-06-10 05:50:32 +08:00
|
|
|
NewASC = ConstantExpr::getAddrSpaceCast(NewGEP, GEP->getType());
|
2014-04-04 05:18:25 +08:00
|
|
|
}
|
2015-06-10 05:50:32 +08:00
|
|
|
return NewASC;
|
2014-04-04 05:18:25 +08:00
|
|
|
}
|
|
|
|
|
2015-06-10 05:50:32 +08:00
|
|
|
Value *NVPTXFavorNonGenericAddrSpaces::hoistAddrSpaceCastFromBitCast(
|
2015-05-30 01:00:27 +08:00
|
|
|
BitCastOperator *BC, int Depth) {
|
2015-06-10 05:50:32 +08:00
|
|
|
Value *NewOperand = hoistAddrSpaceCastFrom(BC->getOperand(0), Depth + 1);
|
|
|
|
if (NewOperand == nullptr)
|
|
|
|
return nullptr;
|
2014-04-04 05:18:25 +08:00
|
|
|
|
2015-06-10 05:50:32 +08:00
|
|
|
// hoistAddrSpaceCastFrom returns an eliminable addrspacecast or nullptr.
|
|
|
|
assert(isEliminableAddrSpaceCast(NewOperand));
|
|
|
|
Operator *Cast = cast<Operator>(NewOperand);
|
2015-05-30 01:00:27 +08:00
|
|
|
|
|
|
|
// Cast = addrspacecast Src
|
|
|
|
// BC = bitcast Cast
|
|
|
|
// =>
|
|
|
|
// Cast' = bitcast Src
|
|
|
|
// BC' = addrspacecast Cast'
|
|
|
|
Value *Src = Cast->getOperand(0);
|
|
|
|
Type *TypeOfNewCast =
|
|
|
|
PointerType::get(BC->getType()->getPointerElementType(),
|
|
|
|
Src->getType()->getPointerAddressSpace());
|
2015-06-10 05:50:32 +08:00
|
|
|
Value *NewBC;
|
2015-05-30 01:00:27 +08:00
|
|
|
if (BitCastInst *BCI = dyn_cast<BitCastInst>(BC)) {
|
|
|
|
Value *NewCast = new BitCastInst(Src, TypeOfNewCast, "", BCI);
|
2015-06-10 05:50:32 +08:00
|
|
|
NewBC = new AddrSpaceCastInst(NewCast, BC->getType(), "", BCI);
|
2015-05-30 01:00:27 +08:00
|
|
|
NewBC->takeName(BC);
|
2015-06-10 05:50:32 +08:00
|
|
|
// Without RAUWing BC, the compiler would visit BC again and emit
|
|
|
|
// redundant instructions. This is exercised in test @rauw in
|
|
|
|
// access-non-generic.ll.
|
2015-05-30 01:00:27 +08:00
|
|
|
BC->replaceAllUsesWith(NewBC);
|
|
|
|
} else {
|
|
|
|
// BC is a constant expression.
|
|
|
|
Constant *NewCast =
|
|
|
|
ConstantExpr::getBitCast(cast<Constant>(Src), TypeOfNewCast);
|
2015-06-10 05:50:32 +08:00
|
|
|
NewBC = ConstantExpr::getAddrSpaceCast(NewCast, BC->getType());
|
2014-04-04 05:18:25 +08:00
|
|
|
}
|
2015-06-10 05:50:32 +08:00
|
|
|
return NewBC;
|
2015-05-30 01:00:27 +08:00
|
|
|
}
|
|
|
|
|
2015-06-10 05:50:32 +08:00
|
|
|
Value *NVPTXFavorNonGenericAddrSpaces::hoistAddrSpaceCastFrom(Value *V,
|
|
|
|
int Depth) {
|
|
|
|
// Returns V if V is already an eliminable addrspacecast.
|
2015-05-30 01:00:27 +08:00
|
|
|
if (isEliminableAddrSpaceCast(V))
|
2015-06-10 05:50:32 +08:00
|
|
|
return V;
|
2015-05-30 01:00:27 +08:00
|
|
|
|
|
|
|
// Limit the depth to prevent this recursive function from running too long.
|
|
|
|
const int MaxDepth = 20;
|
|
|
|
if (Depth >= MaxDepth)
|
2015-06-10 05:50:32 +08:00
|
|
|
return nullptr;
|
2015-05-30 01:00:27 +08:00
|
|
|
|
|
|
|
// If V is a GEP or bitcast, hoist the addrspacecast if any from its pointer
|
|
|
|
// operand. This enables optimizeMemoryInstruction to shortcut addrspacecasts
|
|
|
|
// that are not directly used by the load/store.
|
|
|
|
if (GEPOperator *GEP = dyn_cast<GEPOperator>(V))
|
|
|
|
return hoistAddrSpaceCastFromGEP(GEP, Depth);
|
|
|
|
|
|
|
|
if (BitCastOperator *BC = dyn_cast<BitCastOperator>(V))
|
|
|
|
return hoistAddrSpaceCastFromBitCast(BC, Depth);
|
2014-04-04 05:18:25 +08:00
|
|
|
|
2015-06-10 05:50:32 +08:00
|
|
|
return nullptr;
|
2014-04-04 05:18:25 +08:00
|
|
|
}
|
|
|
|
|
2015-05-30 01:00:27 +08:00
|
|
|
bool NVPTXFavorNonGenericAddrSpaces::optimizeMemoryInstruction(Instruction *MI,
|
|
|
|
unsigned Idx) {
|
2015-06-10 05:50:32 +08:00
|
|
|
Value *NewOperand = hoistAddrSpaceCastFrom(MI->getOperand(Idx));
|
|
|
|
if (NewOperand == nullptr)
|
|
|
|
return false;
|
|
|
|
|
|
|
|
// load/store (addrspacecast X) => load/store X if shortcutting the
|
|
|
|
// addrspacecast is valid and can improve performance.
|
|
|
|
//
|
|
|
|
// e.g.,
|
|
|
|
// %1 = addrspacecast float addrspace(3)* %0 to float*
|
|
|
|
// %2 = load float* %1
|
|
|
|
// ->
|
|
|
|
// %2 = load float addrspace(3)* %0
|
|
|
|
//
|
|
|
|
// Note: the addrspacecast can also be a constant expression.
|
|
|
|
assert(isEliminableAddrSpaceCast(NewOperand));
|
|
|
|
Operator *ASC = dyn_cast<Operator>(NewOperand);
|
|
|
|
MI->setOperand(Idx, ASC->getOperand(0));
|
|
|
|
return true;
|
2015-05-30 01:00:27 +08:00
|
|
|
}
|
|
|
|
|
2014-04-04 05:18:25 +08:00
|
|
|
bool NVPTXFavorNonGenericAddrSpaces::runOnFunction(Function &F) {
|
2016-04-27 07:44:31 +08:00
|
|
|
if (DisableFavorNonGeneric || skipFunction(F))
|
2014-04-04 05:18:25 +08:00
|
|
|
return false;
|
|
|
|
|
|
|
|
bool Changed = false;
|
2015-10-20 08:54:09 +08:00
|
|
|
for (BasicBlock &B : F) {
|
|
|
|
for (Instruction &I : B) {
|
2014-04-04 05:18:25 +08:00
|
|
|
if (isa<LoadInst>(I)) {
|
|
|
|
// V = load P
|
2015-10-20 08:54:09 +08:00
|
|
|
Changed |= optimizeMemoryInstruction(&I, 0);
|
2014-04-04 05:18:25 +08:00
|
|
|
} else if (isa<StoreInst>(I)) {
|
|
|
|
// store V, P
|
2015-10-20 08:54:09 +08:00
|
|
|
Changed |= optimizeMemoryInstruction(&I, 1);
|
2014-04-04 05:18:25 +08:00
|
|
|
}
|
|
|
|
}
|
|
|
|
}
|
|
|
|
return Changed;
|
|
|
|
}
|
|
|
|
|
|
|
|
FunctionPass *llvm::createNVPTXFavorNonGenericAddrSpacesPass() {
|
|
|
|
return new NVPTXFavorNonGenericAddrSpaces();
|
|
|
|
}
|