-//===-- NVPTXTargetTransformInfo.cpp - NVPTX specific TTI pass ---------===//
+//===-- NVPTXTargetTransformInfo.cpp - NVPTX specific TTI -----------------===//
//
// The LLVM Compiler Infrastructure
//
// License. See LICENSE.TXT for details.
//
//===----------------------------------------------------------------------===//
-//
-// \file
-// This file implements a TargetTransformInfo analysis pass specific to the
-// NVPTX target machine. It uses the target's detailed information to provide
-// more precise answers to certain TTI queries, while letting the target
-// independent and default TTI implementations handle the rest.
-//
-//===----------------------------------------------------------------------===//
-#include "NVPTXTargetMachine.h"
+#include "NVPTXTargetTransformInfo.h"
+#include "NVPTXUtilities.h"
#include "llvm/Analysis/LoopInfo.h"
#include "llvm/Analysis/TargetTransformInfo.h"
#include "llvm/Analysis/ValueTracking.h"
#define DEBUG_TYPE "NVPTXtti"
-namespace {
-
-class NVPTXTTIImpl : public BasicTTIImplBase<NVPTXTTIImpl> {
- typedef BasicTTIImplBase<NVPTXTTIImpl> BaseT;
- typedef TargetTransformInfo TTI;
-
- const NVPTXTargetLowering *TLI;
+// Whether the given intrinsic reads threadIdx.x/y/z.
+static bool readsThreadIndex(const IntrinsicInst *II) {
+ switch (II->getIntrinsicID()) {
+ default: return false;
+ case Intrinsic::nvvm_read_ptx_sreg_tid_x:
+ case Intrinsic::nvvm_read_ptx_sreg_tid_y:
+ case Intrinsic::nvvm_read_ptx_sreg_tid_z:
+ return true;
+ }
+}
-public:
- explicit NVPTXTTIImpl(const NVPTXTargetMachine *TM = nullptr)
- : BaseT(TM),
- TLI(TM ? TM->getSubtargetImpl()->getTargetLowering() : nullptr) {}
+static bool readsLaneId(const IntrinsicInst *II) {
+ return II->getIntrinsicID() == Intrinsic::ptx_read_laneid;
+}
- // Provide value semantics. MSVC requires that we spell all of these out.
- NVPTXTTIImpl(const NVPTXTTIImpl &Arg)
- : BaseT(static_cast<const BaseT &>(Arg)), TLI(Arg.TLI) {}
- NVPTXTTIImpl(NVPTXTTIImpl &&Arg)
- : BaseT(std::move(static_cast<BaseT &>(Arg))), TLI(std::move(Arg.TLI)) {}
- NVPTXTTIImpl &operator=(const NVPTXTTIImpl &RHS) {
- BaseT::operator=(static_cast<const BaseT &>(RHS));
- TLI = RHS.TLI;
- return *this;
+// Whether the given intrinsic is an atomic instruction in PTX.
+static bool isNVVMAtomic(const IntrinsicInst *II) {
+ switch (II->getIntrinsicID()) {
+ default: return false;
+ case Intrinsic::nvvm_atomic_load_add_f32:
+ case Intrinsic::nvvm_atomic_load_inc_32:
+ case Intrinsic::nvvm_atomic_load_dec_32:
+ return true;
}
- NVPTXTTIImpl &operator=(NVPTXTTIImpl &&RHS) {
- BaseT::operator=(std::move(static_cast<BaseT &>(RHS)));
- TLI = std::move(RHS.TLI);
- return *this;
- }
-
- bool hasBranchDivergence() { return true; }
+}
- unsigned getArithmeticInstrCost(
- unsigned Opcode, Type *Ty,
- TTI::OperandValueKind Opd1Info = TTI::OK_AnyValue,
- TTI::OperandValueKind Opd2Info = TTI::OK_AnyValue,
- TTI::OperandValueProperties Opd1PropInfo = TTI::OP_None,
- TTI::OperandValueProperties Opd2PropInfo = TTI::OP_None);
-};
+bool NVPTXTTIImpl::isSourceOfDivergence(const Value *V) {
+ // Without inter-procedural analysis, we conservatively assume that arguments
+ // to __device__ functions are divergent.
+ if (const Argument *Arg = dyn_cast<Argument>(V))
+ return !isKernelFunction(*Arg->getParent());
-} // end anonymous namespace
+ if (const Instruction *I = dyn_cast<Instruction>(V)) {
+ // Without pointer analysis, we conservatively assume values loaded from
+ // generic or local address space are divergent.
+ if (const LoadInst *LI = dyn_cast<LoadInst>(I)) {
+ unsigned AS = LI->getPointerAddressSpace();
+ return AS == ADDRESS_SPACE_GENERIC || AS == ADDRESS_SPACE_LOCAL;
+ }
+ // Atomic instructions may cause divergence. Atomic instructions are
+ // executed sequentially across all threads in a warp. Therefore, an earlier
+ // executed thread may see different memory inputs than a later executed
+ // thread. For example, suppose *a = 0 initially.
+ //
+ // atom.global.add.s32 d, [a], 1
+ //
+ // returns 0 for the first thread that enters the critical region, and 1 for
+ // the second thread.
+ if (I->isAtomic())
+ return true;
+ if (const IntrinsicInst *II = dyn_cast<IntrinsicInst>(I)) {
+ // Instructions that read threadIdx are obviously divergent.
+ if (readsThreadIndex(II) || readsLaneId(II))
+ return true;
+ // Handle the NVPTX atomic instrinsics that cannot be represented as an
+ // atomic IR instruction.
+ if (isNVVMAtomic(II))
+ return true;
+ }
+ // Conservatively consider the return value of function calls as divergent.
+ // We could analyze callees with bodies more precisely using
+ // inter-procedural analysis.
+ if (isa<CallInst>(I))
+ return true;
+ }
-ImmutablePass *
-llvm::createNVPTXTargetTransformInfoPass(const NVPTXTargetMachine *TM) {
- return new TargetTransformInfoWrapperPass(NVPTXTTIImpl(TM));
+ return false;
}
-unsigned NVPTXTTIImpl::getArithmeticInstrCost(
+int NVPTXTTIImpl::getArithmeticInstrCost(
unsigned Opcode, Type *Ty, TTI::OperandValueKind Opd1Info,
TTI::OperandValueKind Opd2Info, TTI::OperandValueProperties Opd1PropInfo,
TTI::OperandValueProperties Opd2PropInfo) {
// Legalize the type.
- std::pair<unsigned, MVT> LT = TLI->getTypeLegalizationCost(Ty);
+ std::pair<int, MVT> LT = TLI->getTypeLegalizationCost(DL, Ty);
int ISD = TLI->InstructionOpcodeToISD(Opcode);
Opd1PropInfo, Opd2PropInfo);
}
}
+
+void NVPTXTTIImpl::getUnrollingPreferences(Loop *L,
+ TTI::UnrollingPreferences &UP) {
+ BaseT::getUnrollingPreferences(L, UP);
+
+ // Enable partial unrolling and runtime unrolling, but reduce the
+ // threshold. This partially unrolls small loops which are often
+ // unrolled by the PTX to SASS compiler and unrolling earlier can be
+ // beneficial.
+ UP.Partial = UP.Runtime = true;
+ UP.PartialThreshold = UP.Threshold / 4;
+}
projects / oota-llvm.git / blobdiff