Divergence analysis for GPU programs

[jingyue]

Summary

Some optimizations such as jump threading and loop unswitching can negatively
affect performance when applied to divergent branches. The divergence analysis
added in this patch conservatively estimates which branches in a GPU program
can diverge. This information can then help LLVM to run certain optimizations
selectively.

[jingyue]

+ Tom

[jingyue]

Providing more context, we discussed (http://lists.cs.uiuc.edu/pipermail/llvmdev/2015-January/081126.html) this issue of certain optimizations negatively affecting performance of GPU programs. Per that discussion, the approach of adding automated divergence analysis seems more favorable.

This patch only defines the divergence analysis and does not use it anywhere yet. As you can imagine, after this patch, we can run JumpThreading and LoopUnswitch only on non-divergent branches.

[echristo]

Drive by review: Needs some code comments explaining the functions (though I admit that most of them are pretty straightforward so just a simple "this is the point of this function would be good). Otherwise I don't see anything in particular. Nice big block comment at the top with the explanation. :)

-eric

[• dberlin]

+fernando, who wrote the paper on divergence analysis you cite (and some
follow-ons, i believe), and had an LLVM implementation of it.
It may make sense for him to look over the algorithm itself.

(Note that looking at your paper, i'll point out their is a version that
handles affine constraints, etc at
https://hal.inria.fr/hal-00650235v1/document. I have no opinion on whether
it's better for your usecases, but I know it can be made as fast as your
current algorithm :P)

[• tstellarAMD]

Hi Jingyue,

Thanks for working on this, it looks really interesting. Just one small comment below.

test/Analysis/DivergenceAnalysis/NVPTX/diverge.ll

1–2

Shouldn't this test have a RUN line?

[jingyue]

adds the missing RUN line

[jingyue]

test/Analysis/DivergenceAnalysis/NVPTX/diverge.ll

1–2

Good catch! Believe me I did run them manually :P

[jingyue]

Fernando,

I realized the sync dependency part in my current implementation is
problematic -- I didn't consider the \eta function mentioned in your
papers. I'll fix it and let you know. At the same time, best luck with your
OOPSLA submission!

Jingyue

[madhur13490]

Any plans to improve this further by leveraging AliasAnalysis?

[jingyue]

Hi Madhur,

Absolutely! I mentioned that already in limitations/TODOs. Let me know if you have a specific use case that I could focus on.

Jingyue

[bjarke.roune]

lib/Analysis/DivergenceAnalysis.cpp

177

Can Cond ever be null here, given that we only get to this point if *TI has been marked as a potentially divergent terminator instruction?

180

(ignore - for some reason Phabricator won't let me delete this)

183

More phi nodes than these might need to be marked as diverging if diverging warps can be recognized by the hardware to converge at a point prior to the immediate post-dominator based on the (dynamic) path taken by each diverging subset of threads.

186

It's better to only make one call to getFirstNonPHI(), since it runs in linear time, so this loop is otherwise quadratic in the number of phi nodes:

http://llvm.org/docs/doxygen/html/BasicBlock_8cpp_source.html#l00161

212

(ignore - for some reason Phabricator won't let me delete this)

213

Does any terminator instruction have a value? If not, I think this could be in an else branch.

lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp

55

If all the threads in a warp load the same address at the same time, I think that they should all get the same value. If that's right, then the analysis would remain conservative by letting loads of non-divergent pointers yield non-divergent values, regardless of aliasing.

58

an -> a

68

abviously -> obviously

71

which -> that

[jingyue]

All other comments are ack'ed and will be addressed.

lib/Analysis/DivergenceAnalysis.cpp

177

Since we run reachability on the data flow graph, the Worklist may contain TerminatorInsts (e.g., InvokeInst) that belong to none of the above three types.

183

This is related to the bug I noticed in dealing with values defined in loops. I need to think more about it.

186

Nice catch! Will fix.

213

InvokeInst. Even if not, I'd still prefer doing sync and data dependency in separate stages, just to avoid some mental effort reasoning about "else" :)

lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp

55

That's a very good point! We can be optimistic for at least shared and const loads.

The local address space is private to each thread. Consequently, generic loads need to be treated conservatively, because their addresses may be converted from local addresses.

[bjarke.roune]

lib/Analysis/DivergenceAnalysis.cpp

177

Would a check for >= 2 successors do the trick? If not, is it necessary to pull out the conditions, or would it suffice to have a boolean HasCond? I wonder if BI->getCondition(), IBI->getAdress() and SI->getCondition() can ever be null.

lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp

51

The return value of function calls can also be divergent. I think this function doesn't account for that?

55

I hadn't considered local memory. If we marked local addresses as sources of divergence, the only way I can think of where that might go wrong right now is if a local pointer value is stored into global memory and then read back again, then the thus laundered local pointer is used to store a divergent value and then that divergent value is read back again without being marked as divergent. I wonder if that's a common use case.

[jingyue]

This update fixes sync dependency computation. If a value is used outside of
the loop in that it is defined, the user is sync dependent on the exit
condition of the loop.

lib/Analysis/DivergenceAnalysis.cpp

177

InvokeInst also has two successors. BI->getCondition(), IBI->getAddress() and SI->getCondition() are all non-null.

186

Done.

lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp

51

You're right. Done.

55

Bjarke, I am not sure I understand what you mean by "might go wrong".

Marking values loaded from local address space as divergent is sound, i.e., if our analysis says a branch is non-divergent, it is indeed non-divergent.

58

Done. Thanks!

68

Done.

71

Done.

[jingyue]

Generalize the interface to isDivergent(Value *).

Besides enabling more users of this analysis, this generalizaton makes the
testing easier too. When testing whether a value is divergent, we used to add a
dummy branch that conditions on it. Now we can directly query the uniformity of
this value.

[jingyue]

consider llvm.ptx.read.laneid divergent

[jholewinski]

Looks good to me! Thanks for working on this!

include/llvm/Analysis/Passes.h

namespace llvm {
   //
   // createDelinearizationPass - This pass implements attempts to restore
   // multidimensional array indices from linearized expressions.
   //
   FunctionPass *createDelinearizationPass();
 
   //===--------------------------------------------------------------------===//
   //
+  // createDivergenceAnalysisPass - This pass determines which branches in a GPU
+  // program are divergent.
+  //
+  FunctionPass *createDivergenceAnalysisPass();
+
+  //===--------------------------------------------------------------------===//
+  //
   // Minor pass prototypes, allowing us to expose them through bugpoint and
   // analyze.
   FunctionPass *createInstCountPass();
 
   //===--------------------------------------------------------------------===//
   //
   // createRegionInfoPass - This pass finds all single entry single exit regions
   // in a function and builds the region hierarchy.

include/llvm/Analysis/TargetTransformInfo.h

public:
   /// folded into a load or other instruction, but if they are used in some
   /// other context they may not be folded. This routine can distinguish such
   /// cases.
   ///
   /// The returned cost is defined in terms of \c TargetCostConstants, see its
   /// comments for a detailed explanation of the cost values.
   unsigned getUserCost(const User *U) const;
 
-  /// \brief hasBranchDivergence - Return true if branch divergence exists.
+  /// \brief Return true if branch divergence exists.
+  ///
   /// Branch divergence has a significantly negative impact on GPU performance
   /// when threads in the same wavefront take different paths due to conditional
   /// branches.
   bool hasBranchDivergence() const;
 
+  /// \brief Returns whether V is a source of divergence.
+  ///
+  /// This function provides the target-dependent information for
+  /// the target-independent DivergenceAnalysis. DivergenceAnalysis first
+  /// builds the dependency graph, and then runs the reachability algorithm
+  /// starting with the sources of divergence.
+  bool isSourceOfDivergence(const Value *V) const;
+
   /// \brief Test whether calls to a function lower to actual program function
   /// calls.
   ///
   /// The idea is to test whether the program is likely to require a 'call'
   /// instruction or equivalent in order to call the given function.
   ///
   /// FIXME: It's not clear that this is a good or useful query API. Client's
   /// should probably move to simpler cost metrics using the above.
public:
   virtual unsigned getCallCost(const Function *F,
                                ArrayRef<const Value *> Arguments) = 0;
   virtual unsigned getIntrinsicCost(Intrinsic::ID IID, Type *RetTy,
                                     ArrayRef<Type *> ParamTys) = 0;
   virtual unsigned getIntrinsicCost(Intrinsic::ID IID, Type *RetTy,
                                     ArrayRef<const Value *> Arguments) = 0;
   virtual unsigned getUserCost(const User *U) = 0;
   virtual bool hasBranchDivergence() = 0;
+  virtual bool isSourceOfDivergence(const Value *V) = 0;
   virtual bool isLoweredToCall(const Function *F) = 0;
   virtual void getUnrollingPreferences(Loop *L, UnrollingPreferences &UP) = 0;
   virtual bool isLegalAddImmediate(int64_t Imm) = 0;
   virtual bool isLegalICmpImmediate(int64_t Imm) = 0;
   virtual bool isLegalAddressingMode(Type *Ty, GlobalValue *BaseGV,
                                      int64_t BaseOffset, bool HasBaseReg,
                                      int64_t Scale) = 0;
   virtual bool isLegalMaskedStore(Type *DataType, int Consecutive) = 0;
unsigned getIntrinsicCost(Intrinsic::ID IID, Type *RetTy,
     return Impl.getIntrinsicCost(IID, RetTy, ParamTys);
   }
   unsigned getIntrinsicCost(Intrinsic::ID IID, Type *RetTy,
                             ArrayRef<const Value *> Arguments) override {
     return Impl.getIntrinsicCost(IID, RetTy, Arguments);
   }
   unsigned getUserCost(const User *U) override { return Impl.getUserCost(U); }
   bool hasBranchDivergence() override { return Impl.hasBranchDivergence(); }
+  bool isSourceOfDivergence(const Value *V) override {
+    return Impl.isSourceOfDivergence(V);
+  }
   bool isLoweredToCall(const Function *F) override {
     return Impl.isLoweredToCall(F);
   }
   void getUnrollingPreferences(Loop *L, UnrollingPreferences &UP) override {
     return Impl.getUnrollingPreferences(L, UP);
   }
   bool isLegalAddImmediate(int64_t Imm) override {
     return Impl.isLegalAddImmediate(Imm);

include/llvm/Analysis/TargetTransformInfoImpl.h

unsigned getIntrinsicCost(Intrinsic::ID IID, Type *RetTy,
     case Intrinsic::experimental_gc_relocate:
       // These intrinsics don't actually represent code after lowering.
       return TTI::TCC_Free;
     }
   }
 
   bool hasBranchDivergence() { return false; }
 
+  bool isSourceOfDivergence(const Value *V) { return false; }
+
   bool isLoweredToCall(const Function *F) {
     // FIXME: These should almost certainly not be handled here, and instead
     // handled with the help of TLI or the target itself. This was largely
     // ported from existing analysis heuristics here so that such refactorings
     // can take place in the future.
 
     if (F->isIntrinsic())
       return false;

include/llvm/CodeGen/BasicTTIImpl.h

BasicTTIImplBase &operator=(BasicTTIImplBase &&RHS) {
     return *this;
   }
 
   /// \name Scalar TTI Implementations
   /// @{
 
   bool hasBranchDivergence() { return false; }
 
+  bool isSourceOfDivergence(const Value *V) { return false; }
+
   bool isLegalAddImmediate(int64_t imm) {
     return getTLI()->isLegalAddImmediate(imm);
   }
 
   bool isLegalICmpImmediate(int64_t imm) {
     return getTLI()->isLegalICmpImmediate(imm);
   }
 

include/llvm/InitializePasses.h

 void initializeDAEPass(PassRegistry&);
 void initializeDAHPass(PassRegistry&);
 void initializeDCEPass(PassRegistry&);
 void initializeDSEPass(PassRegistry&);
 void initializeDeadInstEliminationPass(PassRegistry&);
 void initializeDeadMachineInstructionElimPass(PassRegistry&);
 void initializeDelinearizationPass(PassRegistry &);
 void initializeDependenceAnalysisPass(PassRegistry&);
+void initializeDivergenceAnalysisPass(PassRegistry&);
 void initializeDomOnlyPrinterPass(PassRegistry&);
 void initializeDomOnlyViewerPass(PassRegistry&);
 void initializeDomPrinterPass(PassRegistry&);
 void initializeDomViewerPass(PassRegistry&);
 void initializeDominanceFrontierPass(PassRegistry&);
 void initializeDominatorTreeWrapperPassPass(PassRegistry&);
 void initializeEarlyIfConverterPass(PassRegistry&);
 void initializeEdgeBundlesPass(PassRegistry&);

include/llvm/LinkAllPasses.h

ForcePassLinking() {
       (void) llvm::createConstantMergePass();
       (void) llvm::createConstantPropagationPass();
       (void) llvm::createCostModelAnalysisPass();
       (void) llvm::createDeadArgEliminationPass();
       (void) llvm::createDeadCodeEliminationPass();
       (void) llvm::createDeadInstEliminationPass();
       (void) llvm::createDeadStoreEliminationPass();
       (void) llvm::createDependenceAnalysisPass();
+      (void) llvm::createDivergenceAnalysisPass();
       (void) llvm::createDomOnlyPrinterPass();
       (void) llvm::createDomPrinterPass();
       (void) llvm::createDomOnlyViewerPass();
       (void) llvm::createDomViewerPass();
       (void) llvm::createGCOVProfilerPass();
       (void) llvm::createInstrProfilingPass();
       (void) llvm::createFunctionInliningPass();
       (void) llvm::createAlwaysInlinerPass();

lib/Analysis/Analysis.cpp

void llvm::initializeAnalysis(PassRegistry &Registry) {
   initializeCostModelAnalysisPass(Registry);
   initializeCFGViewerPass(Registry);
   initializeCFGPrinterPass(Registry);
   initializeCFGOnlyViewerPass(Registry);
   initializeCFGOnlyPrinterPass(Registry);
   initializeCFLAliasAnalysisPass(Registry);
   initializeDependenceAnalysisPass(Registry);
   initializeDelinearizationPass(Registry);
+  initializeDivergenceAnalysisPass(Registry);
   initializeDominanceFrontierPass(Registry);
   initializeDomViewerPass(Registry);
   initializeDomPrinterPass(Registry);
   initializeDomOnlyViewerPass(Registry);
   initializePostDomViewerPass(Registry);
   initializeDomOnlyPrinterPass(Registry);
   initializePostDomPrinterPass(Registry);
   initializePostDomOnlyViewerPass(Registry);

lib/Analysis/CMakeLists.txt

add_llvm_library(LLVMAnalysis
   CFLAliasAnalysis.cpp
   CGSCCPassManager.cpp
   CaptureTracking.cpp
   CostModel.cpp
   CodeMetrics.cpp
   ConstantFolding.cpp
   Delinearization.cpp
   DependenceAnalysis.cpp
+  DivergenceAnalysis.cpp
   DomPrinter.cpp
   DominanceFrontier.cpp
   IVUsers.cpp
   InstCount.cpp
   InstructionSimplify.cpp
   Interval.cpp
   IntervalPartition.cpp
   LazyCallGraph.cpp

lib/Analysis/DivergenceAnalysis.cpp

+//===- DivergenceAnalysis.cpp ------ Divergence Analysis ------------------===//
+//
+//                     The LLVM Compiler Infrastructure
+//
+// This file is distributed under the University of Illinois Open Source
+// License. See LICENSE.TXT for details.
+//
+//===----------------------------------------------------------------------===//
+//
+// This file defines divergence analysis which determines whether a branch in a
+// GPU program is divergent. It can help branch optimizations such as jump
+// threading and loop unswitching to make better decisions.
+//
+// GPU programs typically use the SIMD execution model, where multiple threads
+// in the same exeuction group have to execute in lock-step. Therefore, if the
+// code contains divergent branches (i.e., threads in a group do not agree on
+// which path of the branch to take), the group of threads has to execute all
+// the paths from that branch with different subsets of threads enabled until
+// they converge at the immediately post-dominating BB of the paths.
+//
+// Due to this execution model, some optimizations such as jump
+// threading and loop unswitching can be unfortunately harmful when performed on
+// divergent branches. Therefore, an analysis that computes which branches in a
+// GPU program are divergent can help the compiler to selectively run these
+// optimizations.
+//
+// This file defines divergence analysis which computes a conservative but
+// non-trivial approximation of all divergent branches in a GPU pogram. It
+// partially implements the approach described in
+//
+//   Divergence Analysis and Optimizations
+//   Coutinho, Sampaio, Pereira, Meira
+//   PACT '11
+//
+// The divergence analysis identifies the sources of divergence (e.g., special
+// variables that hold the thread ID), and recursively marks variables that are
+// data or sync dependent on a source of divergence as divergent.
+//
+// While data dependency is a well-known concept, the notion of sync dependency
+// is worth more explanation. Sync dependence characterizes the control flow
+// aspect of the propagation of branch divergence. For example,
+//
+//   %cond = icmp slt i32 %tid, 10
+//   br i1 %cond, label %then, label %else
+// then:
+//   br label %merge
+// else:
+//   br label %merge
+// merge:
+//   %a = phi i32 [ 0, %then ], [ 1, %else ]
+//
+// Suppose %tid holds the thread ID. Although %a is not data dependent on %tid
+// because %tid is not on its use-def chains, %a is sync dependent on %tid
+// because the branch "br i1 %cond" depends on %tid and affects which value %a
+// is assigned to.
+//
+// Given LLVM's SSA form, computing sync dependency is easy. For any predicate
+// P of a branch, all the PHINodes in the immediate post dominator of the branch
+// are sync dependent on P.
+//
+// The current implementation has the following limitations:
+// 1. intra-procedural. It conservatively considers the arguments of a
+//    non-kernel-entry function as divergent.
+// 2. memory as black box. It conservatively considers the values loaded from
+//    memory as divergent. This can be improved by leveraging pointer analysis.
+//===----------------------------------------------------------------------===//
+
+#include <vector>
+#include "llvm/IR/Dominators.h"
+#include "llvm/ADT/DenseSet.h"
+#include "llvm/Analysis/Passes.h"
+#include "llvm/Analysis/PostDominators.h"
+#include "llvm/Analysis/TargetTransformInfo.h"
+#include "llvm/IR/Function.h"
+#include "llvm/IR/InstIterator.h"
+#include "llvm/IR/Instructions.h"
+#include "llvm/IR/IntrinsicInst.h"
+#include "llvm/IR/Value.h"
+#include "llvm/Pass.h"
+#include "llvm/Support/CommandLine.h"
+#include "llvm/Support/Debug.h"
+#include "llvm/Support/raw_ostream.h"
+using namespace llvm;
+
+#define DEBUG_TYPE "divergence"
+
+namespace {
+class DivergenceAnalysis : public FunctionPass {
+public:
+  static char ID;
+
+  DivergenceAnalysis() : FunctionPass(ID) {
+    initializeDivergenceAnalysisPass(*PassRegistry::getPassRegistry());
+  }
+
+  void getAnalysisUsage(AnalysisUsage &AU) const override {
+    AU.addRequired<PostDominatorTree>();
+    AU.setPreservesAll();
+  }
+
+  bool runOnFunction(Function &F) override;
+
+  void print(raw_ostream &OS, const Module *) const override;
+
+  bool isBranchDivergent(const TerminatorInst *TI) const {
+    return DivergentBranches.count(TI);
+  }
+
+private:
+  DenseSet<const TerminatorInst *> DivergentBranches;
+};
+} // End of anonymous namespace
+
+// Register this pass.
+char DivergenceAnalysis::ID = 0;
+INITIALIZE_PASS_BEGIN(DivergenceAnalysis, "divergence", "Divergence Analysis",
+                      false, true)
+INITIALIZE_PASS_DEPENDENCY(PostDominatorTree)
+INITIALIZE_PASS_END(DivergenceAnalysis, "divergence", "Divergence Analysis",
+                    false, true)
+
+namespace {
+
+class DivergencePropagator {
+public:
+  DivergencePropagator(Function &F, TargetTransformInfo &TTI,
+                       PostDominatorTree &PDT,
+                       DenseSet<const TerminatorInst *> &DB)
+      : F(F), TTI(TTI), PDT(PDT), DB(DB) {}
+  void populateWithSourcesOfDivergence();
+  void propagate();
+
+private:
+  // A helper function that explores data dependents of V.
+  void exploreDataDependency(Value *V);
+  // A helper function that explores sync dependents of V.
+  void exploreSyncDependency(TerminatorInst *TI);
+
+  Function &F;
+  TargetTransformInfo &TTI;
+  PostDominatorTree &PDT;
+  // Place to store the analysis results.
+  DenseSet<const TerminatorInst *> &DB;
+  // Auxiliary data structures for DFS.
+  std::vector<Value *> Worklist;
+  DenseSet<Value *> Visited;
+};
+
+void DivergencePropagator::populateWithSourcesOfDivergence() {
+  Worklist.clear();
+  Visited.clear();
+  for (auto &I : inst_range(F)) {
+    if (TTI.isSourceOfDivergence(&I)) {
+      Worklist.push_back(&I);
+      Visited.insert(&I);
+    }
+  }
+  for (auto &Arg : F.args()) {
+    if (TTI.isSourceOfDivergence(&Arg)) {
+      Worklist.push_back(&Arg);
+      Visited.insert(&Arg);
+    }
+  }
+}
+
+void DivergencePropagator::exploreSyncDependency(TerminatorInst *TI) {
+  Value *Cond = nullptr;
+  if (BranchInst *BI = dyn_cast<BranchInst>(TI)) {
+    if (BI->isConditional())
+      Cond = BI->getCondition();
+  } else if (IndirectBrInst *IBI = dyn_cast<IndirectBrInst>(TI)) {
+    Cond = IBI->getAddress();
+  } else if (SwitchInst *SI = dyn_cast<SwitchInst>(TI)) {
+    Cond = SI->getCondition();
+  }
+  if (Cond == nullptr)
+    return;
+
+  // Since TI is divergent, Cond is also divergent. Per the definition of sync
+  // dependency, we mark all PHINodes in TI's immediate post dominator block as
+  // divergent.
+  BasicBlock *IPostDom = PDT.getNode(TI->getParent())->getIDom()->getBlock();
+  if (IPostDom == nullptr)
+    return;
+  for (auto I = IPostDom->begin(); IPostDom->getFirstNonPHI() != I; ++I) {
+    if (Visited.insert(I).second)
+      Worklist.push_back(I);
+  }
+}
+
+void DivergencePropagator::exploreDataDependency(Value *V) {
+  // Follow def-use chains of V.
+  for (User *U : V->users()) {
+    Instruction *UserInst = cast<Instruction>(U);
+    if (Visited.insert(UserInst).second)
+      Worklist.push_back(UserInst);
+  }
+}
+
+void DivergencePropagator::propagate() {
+  // Traverse the dependency graph using DFS.
+  while (!Worklist.empty()) {
+    Value *V = Worklist.back();
+    Worklist.pop_back();
+    if (TerminatorInst *TI = dyn_cast<TerminatorInst>(V)) {
+      // Terminator instructions with less than two successors are not really
+      // branches, and can be ignored for lookup.
+      if (TI->getNumSuccessors() > 1)
+        DB.insert(TI);
+      exploreSyncDependency(TI);
+    }
+    exploreDataDependency(V);
+  }
+}
+
+} /// end namespace anonymous
+
+FunctionPass *llvm::createDivergenceAnalysisPass() {
+  return new DivergenceAnalysis();
+}
+
+bool DivergenceAnalysis::runOnFunction(Function &F) {
+  auto *TTIWP = getAnalysisIfAvailable<TargetTransformInfoWrapperPass>();
+  if (TTIWP == nullptr)
+    return false;
+
+  TargetTransformInfo &TTI = TTIWP->getTTI(F);
+  // Fast path: if the target does not have branch divergence, we do not mark
+  // any branch as divergent.
+  if (!TTI.hasBranchDivergence())
+    return false;
+
+  DivergentBranches.clear();
+  DivergencePropagator DP(F, TTI, getAnalysis<PostDominatorTree>(),
+                          DivergentBranches);
+  DP.populateWithSourcesOfDivergence();
+  DP.propagate();
+  return false;
+}
+
+void DivergenceAnalysis::print(raw_ostream &OS, const Module *) const {
+  if (DivergentBranches.empty())
+    return;
+  const Function *F = (*DivergentBranches.begin())->getParent()->getParent();
+  // Dumps all divergent branches in F. Iterate instructions using inst_range
+  // to ensure a deterministic order.
+  for (auto &I : inst_range(F)) {
+    if (const TerminatorInst *TI = dyn_cast<TerminatorInst>(&I)) {
+      if (DivergentBranches.count(TI))
+        OS << "DIVERGENT:" << I << "\n";
+    }
+  }
+}

lib/Analysis/TargetTransformInfo.cpp

 unsigned TargetTransformInfo::getUserCost(const User *U) const {
   return TTIImpl->getUserCost(U);
 }
 
 bool TargetTransformInfo::hasBranchDivergence() const {
   return TTIImpl->hasBranchDivergence();
 }
 
+bool TargetTransformInfo::isSourceOfDivergence(const Value *V) const {
+  return TTIImpl->isSourceOfDivergence(V);
+}
+
 bool TargetTransformInfo::isLoweredToCall(const Function *F) const {
   return TTIImpl->isLoweredToCall(F);
 }
 
 void TargetTransformInfo::getUnrollingPreferences(
     Loop *L, UnrollingPreferences &UP) const {
   return TTIImpl->getUnrollingPreferences(L, UP);
 }

lib/Target/NVPTX/NVPTXTargetTransformInfo.h

NVPTXTTIImpl &operator=(NVPTXTTIImpl &&RHS) {
     BaseT::operator=(std::move(static_cast<BaseT &>(RHS)));
     ST = std::move(RHS.ST);
     TLI = std::move(RHS.TLI);
     return *this;
   }
 
   bool hasBranchDivergence() { return true; }
 
+  bool isSourceOfDivergence(const Value *V);
+
   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);
 };
 
 } // end namespace llvm
 
 #endif

lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp

 //===-- NVPTXTargetTransformInfo.cpp - NVPTX specific TTI -----------------===//
 //
 //                     The LLVM Compiler Infrastructure
 //
 // This file is distributed under the University of Illinois Open Source
 // License. See LICENSE.TXT for details.
 //
 //===----------------------------------------------------------------------===//
 
 #include "NVPTXTargetTransformInfo.h"
+#include "NVPTXUtilities.h"
 #include "llvm/Analysis/LoopInfo.h"
 #include "llvm/Analysis/TargetTransformInfo.h"
 #include "llvm/Analysis/ValueTracking.h"
 #include "llvm/CodeGen/BasicTTIImpl.h"
 #include "llvm/Support/Debug.h"
 #include "llvm/Target/CostTable.h"
 #include "llvm/Target/TargetLowering.h"
 using namespace llvm;
 
 #define DEBUG_TYPE "NVPTXtti"
 
+// 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;
+  }
+}
+
+// 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;
+  }
+}
+
+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());
+
+  if (const Instruction *I = dyn_cast<Instruction>(V)) {
+    // Without pointer analysis, we conservatively assume values loaded are
+    // divergent.
+    if (isa<LoadInst>(I))
+      return true;
+    // 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 an 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 abviously divergent.
+      if (readsThreadIndex(II))
+        return true;
+      // Handle the NVPTX atomic instrinsics which cannot be represented as an
+      // atomic IR instruction.
+      if (isNVVMAtomic(II))
+        return true;
+    }
+  }
+
+  return false;
+}
+
 unsigned 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);
 
   int ISD = TLI->InstructionOpcodeToISD(Opcode);

test/Analysis/DivergenceAnalysis/NVPTX/diverge.ll

+target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"
+target triple = "nvptx64-nvidia-cuda"
+
+; return (n < 0 ? a + threadIdx.x : b + threadIdx.x)
+define i32 @no_diverge(i32 %n, i32 %a, i32 %b) {
+; CHECK-LABEL: Printing analysis 'Divergence Analysis' for function 'no_diverge'
+; CHECK-NOT: DIVERGENT
+entry:
+  %tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
+  %cond = icmp slt i32 %n, 0
+  br i1 %cond, label %then, label %else ; not divergent
+then:
+  %a1 = add i32 %a, %tid
+  br label %merge
+else:
+  %b2 = add i32 %b, %tid
+  br label %merge
+merge:
+  %c = phi i32 [ %a1, %then ], [ %b2, %else ]
+  ret i32 %c
+}
+
+; c = a;
+; if (threadIdx.x < 5)    // divergent: data dependent
+;   c = b;
+; return (c == 0 ? a : b) // divergent: sync dependent
+define i32 @sync(i32 %a, i32 %b) {
+; CHECK-LABEL: Printing analysis 'Divergence Analysis' for function 'sync'
+bb1:
+  %tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
+  %cond = icmp slt i32 %tid, 5
+  br i1 %cond, label %bb2, label %bb3
+; CHECK: DIVERGENT: br i1 %cond,
+bb2:
+  br label %bb3
+bb3:
+  %c = phi i32 [ %a, %bb1 ], [ %b, %bb2 ] ; sync dependent on tid
+  %cond2 = icmp eq i32 %c, 0
+  br i1 %cond2, label %bb4, label %bb5 ; therefore divergent
+; CHECK: DIVERGENT: br i1 %cond2,
+bb4:
+  ret i32 %a
+bb5:
+  ret i32 %b
+}
+
+; d = a;
+; if (threadIdx.x >= 5) {  // divergent
+;   c = (n >= 0 ? a : b);  // non-divergent because n is not divergent
+;   d = (c == 0 ? c : b);  // non-divergent because c is not divergent
+; }
+; return (d == 0 ? a : b); // divergent because d is sync dependent on
+;                          // threadIdx.x >= 5
+define i32 @mixed(i32 %n, i32 %a, i32 %b) {
+; CHECK-LABEL: Printing analysis 'Divergence Analysis' for function 'mixed'
+bb1:
+  %tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
+  %cond = icmp slt i32 %tid, 5
+  br i1 %cond, label %bb6, label %bb2
+; CHECK: DIVERGENT: br i1 %cond,
+bb2:
+  %cond2 = icmp slt i32 %n, 0
+  br i1 %cond2, label %bb4, label %bb3
+; CHECK-NOT: DIVERGENT: br i1 %cond2,
+bb3:
+  br label %bb4
+bb4:
+  %c = phi i32 [ %a, %bb2 ], [ %b, %bb3 ]
+  %cond3 = icmp eq i32 %c, 0
+  br i1 %cond3, label %bb6, label %bb5
+; CHECK-NOT: DIVERGENT: br i1 %cond3,
+bb5:
+  br label %bb6
+bb6:
+  %d = phi i32 [ %a, %bb1], [ %b, %bb4 ], [ %c, %bb5 ]
+  %cond4 = icmp eq i32 %d, 0
+  br i1 %cond4, label %bb7, label %bb8
+; CHECK: DIVERGENT: br i1 %cond4,
+bb7:
+  ret i32 %a
+bb8:
+  ret i32 %b
+}
+
+; We conservatively treats all parameters of a __device__ function as divergent.
+define i32 @device(i32 %n, i32 %a, i32 %b) {
+; CHECK-LABEL: Printing analysis 'Divergence Analysis' for function 'device'
+entry:
+  %cond = icmp slt i32 %n, 0
+  br i1 %cond, label %then, label %else
+; CHECK: DIVERGENT: br i1 %cond,
+then:
+  br label %merge
+else:
+  br label %merge
+merge:
+  %c = phi i32 [ %a, %then ], [ %b, %else ]
+  ret i32 %c
+}
+
+declare i32 @llvm.nvvm.read.ptx.sreg.tid.x()
+
+!nvvm.annotations = !{!0, !1, !2}
+!0 = !{i32 (i32, i32, i32)* @no_diverge, !"kernel", i32 1}
+!1 = !{i32 (i32, i32)* @sync, !"kernel", i32 1}
+!2 = !{i32 (i32, i32, i32)* @mixed, !"kernel", i32 1}

test/Analysis/DivergenceAnalysis/NVPTX/lit.local.cfg

+if not 'NVPTX' in config.root.targets:
+    config.unsupported = True