[NewPM] Introduce (GPU)DivergenceAnalysis in the new pass manager

[sameerds]

Summary

The GPUDivergenceAnalysis is now renamed to just "DivergenceAnalysis"
since there is no conflict with LegacyDivergenceAnalysis. In the
legacy PM, this analysis can only be used through the legacy DA
serving as a wrapper. It is now made available as a pass in the new
PM, and has no relation with the legacy DA.

The new DA currently cannot handle irreducible control flow; its
presence can cause the analysis to run indefinitely. The analysis is
now modified to detect this and report all instructions in the
function as divergent. This is super conservative, but allows the
analysis to be used without hanging the compiler.

[aeubanks]

lgtm, thanks for doing this!

[sameerds]

In D96615#2560885, @aeubanks wrote:

lgtm, thanks for doing this!

Thanks, that was fast! I'll wait a couple of days to see if @simoll has any comments about the changes to the analysis implementation itself.

llvm/include/llvm/Analysis/DivergenceAnalysis.h

 class raw_ostream;
 class TargetTransformInfo;
 
 /// \brief Generic divergence analysis for reducible CFGs.
 ///
 /// This analysis propagates divergence in a data-parallel context from sources
 /// of divergence to all users. It requires reducible CFGs. All assignments
 /// should be in SSA form.
-class DivergenceAnalysis {
+class DivergenceAnalysisImpl {
 public:
   /// \brief This instance will analyze the whole function \p F or the loop \p
   /// RegionLoop.
   ///
   /// \param RegionLoop if non-null the analysis is restricted to \p RegionLoop.
   /// Otherwise the whole function is analyzed.
   /// \param IsLCSSAForm whether the analysis may assume that the IR in the
   /// region in in LCSSA form.
-  DivergenceAnalysis(const Function &F, const Loop *RegionLoop,
+  DivergenceAnalysisImpl(const Function &F, const Loop *RegionLoop,
                          const DominatorTree &DT, const LoopInfo &LI,
                          SyncDependenceAnalysis &SDA, bool IsLCSSAForm);
 
   /// \brief The loop that defines the analyzed region (if any).
   const Loop *getRegionLoop() const { return RegionLoop; }
   const Function &getFunction() const { return F; }
 
   /// \brief Whether \p BB is part of the region.
   bool inRegion(const BasicBlock &BB) const;
   /// \brief Whether \p I is part of the region.
public:
 
   /// \brief Whether \p Val is divergent at its definition.
   bool isDivergent(const Value &Val) const;
 
   /// \brief Whether \p U is divergent. Uses of a uniform value can be
   /// divergent.
   bool isDivergentUse(const Use &U) const;
 
-  void print(raw_ostream &OS, const Module *) const;
-
 private:
   /// \brief Mark \p Term as divergent and push all Instructions that become
   /// divergent as a result on the worklist.
   void analyzeControlDivergence(const Instruction &Term);
   /// \brief Mark all phi nodes in \p JoinBlock as divergent and push them on
   /// the worklist.
   void taintAndPushPhiNodes(const BasicBlock &JoinBlock);
 
private:
 
   // Detected/marked divergent values.
   DenseSet<const Value *> DivergentValues;
 
   // Internal worklist for divergence propagation.
   std::vector<const Instruction *> Worklist;
 };
 
-/// \brief Divergence analysis frontend for GPU kernels.
-class GPUDivergenceAnalysis {
-  SyncDependenceAnalysis SDA;
-  DivergenceAnalysis DA;
+class DivergenceInfo {
+  Function &F;
+
+  // If the function contains an irreducible region the divergence
+  // analysis can run indefinitely. We set ContainsIrreducible and no
+  // analysis is actually performed on the function. All values in
+  // this function are conservatively reported as divergent instead.
+  bool ContainsIrreducible;
+  std::unique_ptr<SyncDependenceAnalysis> SDA;
+  std::unique_ptr<DivergenceAnalysisImpl> DA;
 
 public:
-  /// Runs the divergence analysis on @F, a GPU kernel
-  GPUDivergenceAnalysis(Function &F, const DominatorTree &DT,
-                        const PostDominatorTree &PDT, const LoopInfo &LI,
-                        const TargetTransformInfo &TTI);
+  DivergenceInfo(Function &F, const DominatorTree &DT,
+                 const TargetTransformInfo &TTI, bool KnownReducible);
 
   /// Whether any divergence was detected.
-  bool hasDivergence() const { return DA.hasDetectedDivergence(); }
+  bool hasDivergence() const {
+    return ContainsIrreducible || DA->hasDetectedDivergence();
+  }
 
   /// The GPU kernel this analysis result is for
-  const Function &getFunction() const { return DA.getFunction(); }
+  const Function &getFunction() const { return F; }
 
   /// Whether \p V is divergent at its definition.
-  bool isDivergent(const Value &V) const;
+  bool isDivergent(const Value &V) const {
+    return ContainsIrreducible || DA->isDivergent(V);
+  }
 
   /// Whether \p U is divergent. Uses of a uniform value can be divergent.
-  bool isDivergentUse(const Use &U) const;
+  bool isDivergentUse(const Use &U) const {
+    return ContainsIrreducible || DA->isDivergentUse(U);
+  }
 
   /// Whether \p V is uniform/non-divergent.
   bool isUniform(const Value &V) const { return !isDivergent(V); }
 
   /// Whether \p U is uniform/non-divergent. Uses of a uniform value can be
   /// divergent.
   bool isUniformUse(const Use &U) const { return !isDivergentUse(U); }
+};
 
-  /// Print all divergent values in the kernel.
-  void print(raw_ostream &OS, const Module *) const;
+/// \brief Divergence analysis frontend for GPU kernels.
+class DivergenceAnalysis : public AnalysisInfoMixin<DivergenceAnalysis> {
+  friend AnalysisInfoMixin<DivergenceAnalysis>;
+
+  static AnalysisKey Key;
+
+public:
+  using Result = DivergenceInfo;
+
+  /// Runs the divergence analysis on @F, a GPU kernel
+  Result run(Function &F, FunctionAnalysisManager &AM);
 };
 
+/// Printer pass to dump divergence analysis results.
+struct DivergenceAnalysisPrinterPass
+    : public PassInfoMixin<DivergenceAnalysisPrinterPass> {
+  DivergenceAnalysisPrinterPass(raw_ostream &OS) : OS(OS) {}
+
+  PreservedAnalyses run(Function &F, FunctionAnalysisManager &FAM);
+
+private:
+  raw_ostream &OS;
+}; // class DivergenceAnalysisPrinterPass
+
 } // namespace llvm
 
 #endif // LLVM_ANALYSIS_DIVERGENCEANALYSIS_H

llvm/include/llvm/Analysis/LegacyDivergenceAnalysis.h

 #ifndef LLVM_ANALYSIS_LEGACYDIVERGENCEANALYSIS_H
 #define LLVM_ANALYSIS_LEGACYDIVERGENCEANALYSIS_H
 
 #include "llvm/ADT/DenseSet.h"
 #include "llvm/Pass.h"
 #include <memory>
 
 namespace llvm {
+class DivergenceInfo;
 class Function;
-class GPUDivergenceAnalysis;
 class Module;
 class raw_ostream;
 class TargetTransformInfo;
 class Use;
 class Value;
 
 class LegacyDivergenceAnalysis : public FunctionPass {
 public:
public:
   void removeValue(const Value *V) { DivergentValues.erase(V); }
 
 private:
   // Whether analysis should be performed by GPUDivergenceAnalysis.
   bool shouldUseGPUDivergenceAnalysis(const Function &F,
                                       const TargetTransformInfo &TTI) const;
 
   // (optional) handle to new DivergenceAnalysis
-  std::unique_ptr<GPUDivergenceAnalysis> gpuDA;
+  std::unique_ptr<DivergenceInfo> gpuDA;
 
   // Stores all divergent values.
   DenseSet<const Value *> DivergentValues;
 
   // Stores divergent uses of possibly uniform values.
   DenseSet<const Use *> DivergentUses;
 };
 } // End llvm namespace
 
 #endif // LLVM_ANALYSIS_LEGACYDIVERGENCEANALYSIS_H

llvm/lib/Analysis/DivergenceAnalysis.cpp

 // This implementation is derived from the Vectorization Analysis of the
 // Region Vectorizer (RV). That implementation in turn is based on the approach
 // described in
 //
 //   Improving Performance of OpenCL on CPUs
 //   Ralf Karrenberg and Sebastian Hack
 //   CC '12
 //
-// This DivergenceAnalysis implementation is generic in the sense that it does
+// This implementation is generic in the sense that it does
 // not itself identify original sources of divergence.
 // Instead specialized adapter classes, (LoopDivergenceAnalysis) for loops and
-// (GPUDivergenceAnalysis) for GPU programs, identify the sources of divergence
+// (DivergenceAnalysis) for functions, identify the sources of divergence
 // (e.g., special variables that hold the thread ID or the iteration variable).
 //
 // The generic implementation propagates divergence to variables that are data
 // or sync dependent on a source of divergence.
 //
 // 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,

 // 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.
 //
 // The sync dependence detection (which branch induces divergence in which join
 // points) is implemented in the SyncDependenceAnalysis.
 //
-// The current DivergenceAnalysis implementation has the following limitations:
+// The current implementation has the following limitations:
 // 1. intra-procedural. It conservatively considers the arguments of a
 //    non-kernel-entry function and the return value of a function call as
 //    divergent.
 // 2. memory as black box. It conservatively considers values loaded from
 //    generic or local address as divergent. This can be improved by leveraging
 //    pointer analysis and/or by modelling non-escaping memory objects in SSA
 //    as done in RV.
 //
 //===----------------------------------------------------------------------===//
 
 #include "llvm/Analysis/DivergenceAnalysis.h"
+#include "llvm/Analysis/CFG.h"
 #include "llvm/Analysis/LoopInfo.h"
 #include "llvm/Analysis/Passes.h"
 #include "llvm/Analysis/PostDominators.h"
 #include "llvm/Analysis/TargetTransformInfo.h"
 #include "llvm/IR/Dominators.h"
 #include "llvm/IR/InstIterator.h"
 #include "llvm/IR/Instructions.h"
 #include "llvm/IR/IntrinsicInst.h"
 #include "llvm/IR/Value.h"
 #include "llvm/Support/Debug.h"
 #include "llvm/Support/raw_ostream.h"
 
 using namespace llvm;
 
-#define DEBUG_TYPE "divergence-analysis"
+#define DEBUG_TYPE "divergence"
 
-// class DivergenceAnalysis
-DivergenceAnalysis::DivergenceAnalysis(
+DivergenceAnalysisImpl::DivergenceAnalysisImpl(
     const Function &F, const Loop *RegionLoop, const DominatorTree &DT,
     const LoopInfo &LI, SyncDependenceAnalysis &SDA, bool IsLCSSAForm)
     : F(F), RegionLoop(RegionLoop), DT(DT), LI(LI), SDA(SDA),
       IsLCSSAForm(IsLCSSAForm) {}
 
-bool DivergenceAnalysis::markDivergent(const Value &DivVal) {
+bool DivergenceAnalysisImpl::markDivergent(const Value &DivVal) {
   if (isAlwaysUniform(DivVal))
     return false;
   assert(isa<Instruction>(DivVal) || isa<Argument>(DivVal));
   assert(!isAlwaysUniform(DivVal) && "cannot be a divergent");
   return DivergentValues.insert(&DivVal).second;
 }
 
-void DivergenceAnalysis::addUniformOverride(const Value &UniVal) {
+void DivergenceAnalysisImpl::addUniformOverride(const Value &UniVal) {
   UniformOverrides.insert(&UniVal);
 }
 
-bool DivergenceAnalysis::isTemporalDivergent(const BasicBlock &ObservingBlock,
-                                             const Value &Val) const {
+bool DivergenceAnalysisImpl::isTemporalDivergent(
+    const BasicBlock &ObservingBlock, const Value &Val) const {
   const auto *Inst = dyn_cast<const Instruction>(&Val);
   if (!Inst)
     return false;
   // check whether any divergent loop carrying Val terminates before control
   // proceeds to ObservingBlock
   for (const auto *Loop = LI.getLoopFor(Inst->getParent());
        Loop != RegionLoop && !Loop->contains(&ObservingBlock);
        Loop = Loop->getParentLoop()) {
     if (DivergentLoops.contains(Loop))
       return true;
   }
 
   return false;
 }
 
-bool DivergenceAnalysis::inRegion(const Instruction &I) const {
+bool DivergenceAnalysisImpl::inRegion(const Instruction &I) const {
   return I.getParent() && inRegion(*I.getParent());
 }
 
-bool DivergenceAnalysis::inRegion(const BasicBlock &BB) const {
+bool DivergenceAnalysisImpl::inRegion(const BasicBlock &BB) const {
   return (!RegionLoop && BB.getParent() == &F) || RegionLoop->contains(&BB);
 }
 
-void DivergenceAnalysis::pushUsers(const Value &V) {
+void DivergenceAnalysisImpl::pushUsers(const Value &V) {
   const auto *I = dyn_cast<const Instruction>(&V);
 
   if (I && I->isTerminator()) {
     analyzeControlDivergence(*I);
     return;
   }
 
   for (const auto *User : V.users()) {
static const Instruction *getIfCarriedInstruction(const Use &U,
   const auto *I = dyn_cast<const Instruction>(&U);
   if (!I)
     return nullptr;
   if (!DivLoop.contains(I))
     return nullptr;
   return I;
 }
 
-void DivergenceAnalysis::analyzeTemporalDivergence(const Instruction &I,
-                                                   const Loop &OuterDivLoop) {
+void DivergenceAnalysisImpl::analyzeTemporalDivergence(
+    const Instruction &I, const Loop &OuterDivLoop) {
   if (isAlwaysUniform(I))
     return;
   if (isDivergent(I))
     return;
 
   LLVM_DEBUG(dbgs() << "Analyze temporal divergence: " << I.getName() << "\n");
   assert((isa<PHINode>(I) || !IsLCSSAForm) &&
          "In LCSSA form all users of loop-exiting defs are Phi nodes.");
   for (const Use &Op : I.operands()) {
     const auto *OpInst = getIfCarriedInstruction(Op, OuterDivLoop);
     if (!OpInst)
       continue;
     if (markDivergent(I))
       pushUsers(I);
     return;
   }
 }
 
 // marks all users of loop-carried values of the loop headed by LoopHeader as
 // divergent
-void DivergenceAnalysis::analyzeLoopExitDivergence(const BasicBlock &DivExit,
-                                                   const Loop &OuterDivLoop) {
+void DivergenceAnalysisImpl::analyzeLoopExitDivergence(
+    const BasicBlock &DivExit, const Loop &OuterDivLoop) {
   // All users are in immediate exit blocks
   if (IsLCSSAForm) {
     for (const auto &Phi : DivExit.phis()) {
       analyzeTemporalDivergence(Phi, OuterDivLoop);
     }
     return;
   }
 
for (auto *SuccBlock : successors(UserBlock)) {
       if (!Visited.insert(SuccBlock).second) {
         continue;
       }
       TaintStack.push_back(SuccBlock);
     }
   } while (!TaintStack.empty());
 }
 
-void DivergenceAnalysis::propagateLoopExitDivergence(const BasicBlock &DivExit,
-                                                     const Loop &InnerDivLoop) {
+void DivergenceAnalysisImpl::propagateLoopExitDivergence(
+    const BasicBlock &DivExit, const Loop &InnerDivLoop) {
   LLVM_DEBUG(dbgs() << "\tpropLoopExitDiv " << DivExit.getName() << "\n");
 
   // Find outer-most loop that does not contain \p DivExit
   const Loop *DivLoop = &InnerDivLoop;
   const Loop *OuterDivLoop = DivLoop;
   const Loop *ExitLevelLoop = LI.getLoopFor(&DivExit);
   const unsigned LoopExitDepth =
       ExitLevelLoop ? ExitLevelLoop->getLoopDepth() : 0;
   while (DivLoop && DivLoop->getLoopDepth() > LoopExitDepth) {
     DivergentLoops.insert(DivLoop); // all crossed loops are divergent
     OuterDivLoop = DivLoop;
     DivLoop = DivLoop->getParentLoop();
   }
   LLVM_DEBUG(dbgs() << "\tOuter-most left loop: " << OuterDivLoop->getName()
                     << "\n");
 
   analyzeLoopExitDivergence(DivExit, *OuterDivLoop);
 }
 
 // this is a divergent join point - mark all phi nodes as divergent and push
 // them onto the stack.
-void DivergenceAnalysis::taintAndPushPhiNodes(const BasicBlock &JoinBlock) {
+void DivergenceAnalysisImpl::taintAndPushPhiNodes(const BasicBlock &JoinBlock) {
   LLVM_DEBUG(dbgs() << "taintAndPushPhiNodes in " << JoinBlock.getName()
                     << "\n");
 
   // ignore divergence outside the region
   if (!inRegion(JoinBlock)) {
     return;
   }
 
   // push non-divergent phi nodes in JoinBlock to the worklist
   for (const auto &Phi : JoinBlock.phis()) {
     if (isDivergent(Phi))
       continue;
     // FIXME Theoretically ,the 'undef' value could be replaced by any other
     // value causing spurious divergence.
     if (Phi.hasConstantOrUndefValue())
       continue;
     if (markDivergent(Phi))
       Worklist.push_back(&Phi);
   }
 }
 
-void DivergenceAnalysis::analyzeControlDivergence(const Instruction &Term) {
+void DivergenceAnalysisImpl::analyzeControlDivergence(const Instruction &Term) {
   LLVM_DEBUG(dbgs() << "analyzeControlDiv " << Term.getParent()->getName()
                     << "\n");
 
   // Don't propagate divergence from unreachable blocks.
   if (!DT.isReachableFromEntry(Term.getParent()))
     return;
 
   const auto *BranchLoop = LI.getLoopFor(Term.getParent());
 
   const auto &DivDesc = SDA.getJoinBlocks(Term);
 
   // Iterate over all blocks now reachable by a disjoint path join
   for (const auto *JoinBlock : DivDesc.JoinDivBlocks) {
     taintAndPushPhiNodes(*JoinBlock);
   }
 
   assert(DivDesc.LoopDivBlocks.empty() || BranchLoop);
   for (const auto *DivExitBlock : DivDesc.LoopDivBlocks) {
     propagateLoopExitDivergence(*DivExitBlock, *BranchLoop);
   }
 }
 
-void DivergenceAnalysis::compute() {
+void DivergenceAnalysisImpl::compute() {
   // Initialize worklist.
   auto DivValuesCopy = DivergentValues;
   for (const auto *DivVal : DivValuesCopy) {
     assert(isDivergent(*DivVal) && "Worklist invariant violated!");
     pushUsers(*DivVal);
   }
 
   // All values on the Worklist are divergent.
   // Their users may not have been updated yed.
   while (!Worklist.empty()) {
     const Instruction &I = *Worklist.back();
     Worklist.pop_back();
 
     // propagate value divergence to users
     assert(isDivergent(I) && "Worklist invariant violated!");
     pushUsers(I);
   }
 }
 
-bool DivergenceAnalysis::isAlwaysUniform(const Value &V) const {
+bool DivergenceAnalysisImpl::isAlwaysUniform(const Value &V) const {
   return UniformOverrides.contains(&V);
 }
 
-bool DivergenceAnalysis::isDivergent(const Value &V) const {
+bool DivergenceAnalysisImpl::isDivergent(const Value &V) const {
   return DivergentValues.contains(&V);
 }
 
-bool DivergenceAnalysis::isDivergentUse(const Use &U) const {
+bool DivergenceAnalysisImpl::isDivergentUse(const Use &U) const {
   Value &V = *U.get();
   Instruction &I = *cast<Instruction>(U.getUser());
   return isDivergent(V) || isTemporalDivergent(*I.getParent(), V);
 }
 
-void DivergenceAnalysis::print(raw_ostream &OS, const Module *) const {
-  if (DivergentValues.empty())
+DivergenceInfo::DivergenceInfo(Function &F, const DominatorTree &DT,
+                               const PostDominatorTree &PDT, const LoopInfo &LI,
+                               const TargetTransformInfo &TTI,
+                               bool KnownReducible)
+    : F(F), ContainsIrreducible(false) {
+  if (!KnownReducible) {
+    using RPOTraversal = ReversePostOrderTraversal<const Function *>;
+    RPOTraversal FuncRPOT(&F);
+    if (containsIrreducibleCFG<const BasicBlock *, const RPOTraversal,
+                               const LoopInfo>(FuncRPOT, LI)) {
+      ContainsIrreducible = true;
       return;
-  // iterate instructions using instructions() to ensure a deterministic order.
-  for (auto &I : instructions(F)) {
-    if (isDivergent(I))
-      OS << "DIVERGENT:" << I << '\n';
     }
   }
-
-// class GPUDivergenceAnalysis
-GPUDivergenceAnalysis::GPUDivergenceAnalysis(Function &F,
-                                             const DominatorTree &DT,
-                                             const PostDominatorTree &PDT,
-                                             const LoopInfo &LI,
-                                             const TargetTransformInfo &TTI)
-    : SDA(DT, PDT, LI), DA(F, nullptr, DT, LI, SDA, /* LCSSA */ false) {
+  SDA = std::make_unique<SyncDependenceAnalysis>(DT, PDT, LI);
+  DA = std::make_unique<DivergenceAnalysisImpl>(F, nullptr, DT, LI, *SDA,
+                                                /* LCSSA */ false);
   for (auto &I : instructions(F)) {
     if (TTI.isSourceOfDivergence(&I)) {
-      DA.markDivergent(I);
+      DA->markDivergent(I);
     } else if (TTI.isAlwaysUniform(&I)) {
-      DA.addUniformOverride(I);
+      DA->addUniformOverride(I);
     }
   }
   for (auto &Arg : F.args()) {
     if (TTI.isSourceOfDivergence(&Arg)) {
-      DA.markDivergent(Arg);
+      DA->markDivergent(Arg);
     }
   }
 
-  DA.compute();
+  DA->compute();
 }
 
-bool GPUDivergenceAnalysis::isDivergent(const Value &val) const {
-  return DA.isDivergent(val);
-}
+AnalysisKey DivergenceAnalysis::Key;
+
+DivergenceAnalysis::Result
+DivergenceAnalysis::run(Function &F, FunctionAnalysisManager &AM) {
+  auto &DT = AM.getResult<DominatorTreeAnalysis>(F);
+  auto &PDT = AM.getResult<PostDominatorTreeAnalysis>(F);
+  auto &LI = AM.getResult<LoopAnalysis>(F);
+  auto &TTI = AM.getResult<TargetIRAnalysis>(F);
 
-bool GPUDivergenceAnalysis::isDivergentUse(const Use &use) const {
-  return DA.isDivergentUse(use);
+  return DivergenceInfo(F, DT, PDT, LI, TTI, /* KnownReducible = */ false);
 }
 
-void GPUDivergenceAnalysis::print(raw_ostream &OS, const Module *mod) const {
-  OS << "Divergence of kernel " << DA.getFunction().getName() << " {\n";
-  DA.print(OS, mod);
-  OS << "}\n";
+PreservedAnalyses
+DivergenceAnalysisPrinterPass::run(Function &F, FunctionAnalysisManager &FAM) {
+  auto &DI = FAM.getResult<DivergenceAnalysis>(F);
+  OS << "'Divergence Analysis' for function '" << F.getName() << "':\n";
+  if (DI.hasDivergence()) {
+    for (auto &Arg : F.args()) {
+      OS << (DI.isDivergent(Arg) ? "DIVERGENT: " : "           ");
+      OS << Arg << "\n";
+    }
+    for (auto BI = F.begin(), BE = F.end(); BI != BE; ++BI) {
+      auto &BB = *BI;
+      OS << "\n           " << BB.getName() << ":\n";
+      for (auto &I : BB.instructionsWithoutDebug()) {
+        OS << (DI.isDivergent(I) ? "DIVERGENT:     " : "               ");
+        OS << I << "\n";
+      }
+    }
+  }
+  return PreservedAnalyses::all();
 }

llvm/lib/Analysis/LegacyDivergenceAnalysis.cpp

bool LegacyDivergenceAnalysis::runOnFunction(Function &F) {
   gpuDA = nullptr;
 
   auto &DT = getAnalysis<DominatorTreeWrapperPass>().getDomTree();
   auto &PDT = getAnalysis<PostDominatorTreeWrapperPass>().getPostDomTree();
 
   if (shouldUseGPUDivergenceAnalysis(F, TTI)) {
     // run the new GPU divergence analysis
     auto &LI = getAnalysis<LoopInfoWrapperPass>().getLoopInfo();
-    gpuDA = std::make_unique<GPUDivergenceAnalysis>(F, DT, PDT, LI, TTI);
+    gpuDA = std::make_unique<DivergenceInfo>(F, DT, PDT, LI, TTI,
+                                             /* KnownReducible  = */ true);
 
   } else {
     // run LLVM's existing DivergenceAnalysis
     DivergencePropagator DP(F, TTI, DT, PDT, DivergentValues, DivergentUses);
     DP.populateWithSourcesOfDivergence();
     DP.propagate();
   }
 

llvm/lib/Passes/PassBuilder.cpp

 #include "llvm/Analysis/CFLSteensAliasAnalysis.h"
 #include "llvm/Analysis/CGSCCPassManager.h"
 #include "llvm/Analysis/CallGraph.h"
 #include "llvm/Analysis/DDG.h"
 #include "llvm/Analysis/DDGPrinter.h"
 #include "llvm/Analysis/Delinearization.h"
 #include "llvm/Analysis/DemandedBits.h"
 #include "llvm/Analysis/DependenceAnalysis.h"
+#include "llvm/Analysis/DivergenceAnalysis.h"
 #include "llvm/Analysis/DominanceFrontier.h"
 #include "llvm/Analysis/FunctionPropertiesAnalysis.h"
 #include "llvm/Analysis/GlobalsModRef.h"
 #include "llvm/Analysis/IRSimilarityIdentifier.h"
 #include "llvm/Analysis/IVUsers.h"
 #include "llvm/Analysis/InlineAdvisor.h"
 #include "llvm/Analysis/InlineSizeEstimatorAnalysis.h"
 #include "llvm/Analysis/InstCount.h"

llvm/lib/Passes/PassRegistry.def

 FUNCTION_ANALYSIS("opt-remark-emit", OptimizationRemarkEmitterAnalysis())
 FUNCTION_ANALYSIS("scalar-evolution", ScalarEvolutionAnalysis())
 FUNCTION_ANALYSIS("stack-safety-local", StackSafetyAnalysis())
 FUNCTION_ANALYSIS("targetlibinfo", TargetLibraryAnalysis())
 FUNCTION_ANALYSIS("targetir",
                   TM ? TM->getTargetIRAnalysis() : TargetIRAnalysis())
 FUNCTION_ANALYSIS("verify", VerifierAnalysis())
 FUNCTION_ANALYSIS("pass-instrumentation", PassInstrumentationAnalysis(PIC))
+FUNCTION_ANALYSIS("divergence", DivergenceAnalysis())
 
 #ifndef FUNCTION_ALIAS_ANALYSIS
 #define FUNCTION_ALIAS_ANALYSIS(NAME, CREATE_PASS)                             \
   FUNCTION_ANALYSIS(NAME, CREATE_PASS)
 #endif
 FUNCTION_ALIAS_ANALYSIS("basic-aa", BasicAA())
 FUNCTION_ALIAS_ANALYSIS("cfl-anders-aa", CFLAndersAA())
 FUNCTION_ALIAS_ANALYSIS("cfl-steens-aa", CFLSteensAA())

 FUNCTION_PASS("objc-arc-contract", ObjCARCContractPass())
 FUNCTION_PASS("objc-arc-expand", ObjCARCExpandPass())
 FUNCTION_PASS("pgo-memop-opt", PGOMemOPSizeOpt())
 FUNCTION_PASS("print", PrintFunctionPass(dbgs()))
 FUNCTION_PASS("print<assumptions>", AssumptionPrinterPass(dbgs()))
 FUNCTION_PASS("print<block-freq>", BlockFrequencyPrinterPass(dbgs()))
 FUNCTION_PASS("print<branch-prob>", BranchProbabilityPrinterPass(dbgs()))
 FUNCTION_PASS("print<da>", DependenceAnalysisPrinterPass(dbgs()))
+FUNCTION_PASS("print<divergence>", DivergenceAnalysisPrinterPass(dbgs()))
 FUNCTION_PASS("print<domtree>", DominatorTreePrinterPass(dbgs()))
 FUNCTION_PASS("print<postdomtree>", PostDominatorTreePrinterPass(dbgs()))
 FUNCTION_PASS("print<delinearization>", DelinearizationPrinterPass(dbgs()))
 FUNCTION_PASS("print<demanded-bits>", DemandedBitsPrinterPass(dbgs()))
 FUNCTION_PASS("print<domfrontier>", DominanceFrontierPrinterPass(dbgs()))
 FUNCTION_PASS("print<func-properties>", FunctionPropertiesPrinterPass(dbgs()))
 FUNCTION_PASS("print<inline-cost>", InlineCostAnnotationPrinterPass(dbgs()))
 FUNCTION_PASS("print<inliner-size-estimator>", 

llvm/test/Analysis/DivergenceAnalysis/AMDGPU/always_uniform.ll

-; RUN: opt  -mtriple amdgcn-unknown-amdhsa -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt  -mtriple amdgcn-unknown-amdhsa -enable-new-pm=0 -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt -mtriple amdgcn-unknown-amdhsa -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 
 ; CHECK-LABEL: for function 'readfirstlane':
 define amdgpu_kernel void @readfirstlane() {
   %id.x = call i32 @llvm.amdgcn.workitem.id.x()
 ; CHECK: DIVERGENT:  %id.x = call i32 @llvm.amdgcn.workitem.id.x()
   %first.lane = call i32 @llvm.amdgcn.readfirstlane(i32 %id.x)
 ; CHECK-NOT: DIVERGENT:  %first.lane = call i32 @llvm.amdgcn.readfirstlane(i32 %id.x)
   ret void

 ; CHECK-LABEL: for function 'asm_sgpr':
 ; CHECK: DIVERGENT: i32 %divergent
 ; CHECK-NOT: DIVERGENT
 define i32 @asm_sgpr(i32 %divergent) {
   %sgpr = call i32 asm "; def $0, $1","=s,v"(i32 %divergent)
   ret i32 %sgpr
 }
 
-; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'asm_mixed_sgpr_vgpr':
+; CHECK-LABEL: Divergence Analysis' for function 'asm_mixed_sgpr_vgpr':
 ; CHECK: DIVERGENT: %asm = call { i32, i32 } asm "; def $0, $1, $2", "=s,=v,v"(i32 %divergent)
 ; CHECK-NEXT: {{^[ \t]+}}%sgpr = extractvalue { i32, i32 } %asm, 0
 ; CHECK-NEXT: DIVERGENT:       %vgpr = extractvalue { i32, i32 } %asm, 1
 define void @asm_mixed_sgpr_vgpr(i32 %divergent) {
   %asm = call { i32, i32 } asm "; def $0, $1, $2","=s,=v,v"(i32 %divergent)
   %sgpr = extractvalue { i32, i32 } %asm, 0
   %vgpr = extractvalue { i32, i32 } %asm, 1
   store i32 %sgpr, i32 addrspace(1)* undef

llvm/test/Analysis/DivergenceAnalysis/AMDGPU/atomics.ll

-; RUN: opt -mtriple=amdgcn-- -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt -mtriple=amdgcn-- -enable-new-pm=0 -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt -mtriple amdgcn-- -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 
 ; CHECK: DIVERGENT: %orig = atomicrmw xchg i32* %ptr, i32 %val seq_cst
 define i32 @test1(i32* %ptr, i32 %val) #0 {
   %orig = atomicrmw xchg i32* %ptr, i32 %val seq_cst
   ret i32 %orig
 }
 
 ; CHECK: DIVERGENT: %orig = cmpxchg i32* %ptr, i32 %cmp, i32 %new seq_cst seq_cst

llvm/test/Analysis/DivergenceAnalysis/AMDGPU/b42473-r1-crash.ll

-; RUN: opt -mtriple amdgcn-unknown-amdhsa -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt -mtriple amdgcn-unknown-amdhsa -enable-new-pm=0 -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt -mtriple amdgcn-unknown-amdhsa -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 
 declare i32 @gf2(i32)
 declare i32 @gf1(i32)
 
 define  void @tw1(i32 addrspace(4)* noalias nocapture readonly %A, i32 addrspace(4)* noalias nocapture %B) local_unnamed_addr #2 {
-; CHECK: Printing analysis 'Legacy Divergence Analysis' for function 'tw1':
+; CHECK: Divergence Analysis' for function 'tw1':
 ; CHECK: DIVERGENT: i32 addrspace(4)* %A
 ; CHECK: DIVERGENT: i32 addrspace(4)* %B
 entry:
 ; CHECK: DIVERGENT:       %call = tail call i32 @gf2(i32 0) #0
 ; CHECK: DIVERGENT:       %cmp = icmp ult i32 %call, 16
 ; CHECK: DIVERGENT:       br i1 %cmp, label %if.then, label %new_exit
   %call = tail call  i32 @gf2(i32 0) #3
   %cmp = icmp ult i32 %call, 16

llvm/test/Analysis/DivergenceAnalysis/AMDGPU/control-flow-intrinsics.ll

-; RUN: opt -mtriple=amdgcn-mesa-mesa3d -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt -mtriple=amdgcn-mesa-mesa3d -enable-new-pm=0 -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt -mtriple amdgcn-mesa-mesa3d -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 
 ; Tests control flow intrinsics that should be treated as uniform
 
-; CHECK: Printing analysis 'Legacy Divergence Analysis' for function 'test_if_break':
+; CHECK: Divergence Analysis' for function 'test_if_break':
 ; CHECK: DIVERGENT: %cond = icmp eq i32 %arg0, 0
 ; CHECK-NOT: DIVERGENT
 ; CHECK: ret void
 define amdgpu_ps void @test_if_break(i32 %arg0, i64 inreg %saved) {
 entry:
   %cond = icmp eq i32 %arg0, 0
   %break = call i64 @llvm.amdgcn.if.break.i64.i64(i1 %cond, i64 %saved)
   store volatile i64 %break, i64 addrspace(1)* undef
   ret void
 }
 
-; CHECK: Printing analysis 'Legacy Divergence Analysis' for function 'test_if':
+; CHECK: Divergence Analysis' for function 'test_if':
 ; CHECK: DIVERGENT: %cond = icmp eq i32 %arg0, 0
 ; CHECK-NEXT: DIVERGENT: %if = call { i1, i64 } @llvm.amdgcn.if.i64(i1 %cond)
 ; CHECK-NEXT: DIVERGENT: %if.bool = extractvalue { i1, i64 } %if, 0
 ; CHECK-NOT: DIVERGENT
 ; CHECK: DIVERGENT: %if.bool.ext = zext i1 %if.bool to i32
 define void @test_if(i32 %arg0) {
 entry:
   %cond = icmp eq i32 %arg0, 0
   %if = call { i1, i64 } @llvm.amdgcn.if.i64(i1 %cond)
   %if.bool = extractvalue { i1, i64 } %if, 0
   %if.mask = extractvalue { i1, i64 } %if, 1
   %if.bool.ext = zext i1 %if.bool to i32
   store volatile i32 %if.bool.ext, i32 addrspace(1)* undef
   store volatile i64 %if.mask, i64 addrspace(1)* undef
   ret void
 }
 
 ; The result should still be treated as divergent, even with a uniform source.
-; CHECK: Printing analysis 'Legacy Divergence Analysis' for function 'test_if_uniform':
+; CHECK: Divergence Analysis' for function 'test_if_uniform':
 ; CHECK-NOT: DIVERGENT
 ; CHECK: DIVERGENT: %if = call { i1, i64 } @llvm.amdgcn.if.i64(i1 %cond)
 ; CHECK-NEXT: DIVERGENT: %if.bool = extractvalue { i1, i64 } %if, 0
 ; CHECK-NOT: DIVERGENT
 ; CHECK: DIVERGENT: %if.bool.ext = zext i1 %if.bool to i32
 define amdgpu_ps void @test_if_uniform(i32 inreg %arg0) {
 entry:
   %cond = icmp eq i32 %arg0, 0
   %if = call { i1, i64 } @llvm.amdgcn.if.i64(i1 %cond)
   %if.bool = extractvalue { i1, i64 } %if, 0
   %if.mask = extractvalue { i1, i64 } %if, 1
   %if.bool.ext = zext i1 %if.bool to i32
   store volatile i32 %if.bool.ext, i32 addrspace(1)* undef
   store volatile i64 %if.mask, i64 addrspace(1)* undef
   ret void
 }
 
-; CHECK: Printing analysis 'Legacy Divergence Analysis' for function 'test_loop_uniform':
+; CHECK: Divergence Analysis' for function 'test_loop_uniform':
 ; CHECK: DIVERGENT: %loop = call i1 @llvm.amdgcn.loop.i64(i64 %mask)
 define amdgpu_ps void @test_loop_uniform(i64 inreg %mask) {
 entry:
   %loop = call i1 @llvm.amdgcn.loop.i64(i64 %mask)
   %loop.ext = zext i1 %loop to i32
   store volatile i32 %loop.ext, i32 addrspace(1)* undef
   ret void
 }
 
-; CHECK: Printing analysis 'Legacy Divergence Analysis' for function 'test_else':
+; CHECK: Divergence Analysis' for function 'test_else':
 ; CHECK: DIVERGENT: %else = call { i1, i64 } @llvm.amdgcn.else.i64.i64(i64 %mask)
 ; CHECK: DIVERGENT:       %else.bool = extractvalue { i1, i64 } %else, 0
 ; CHECK: {{^[ \t]+}}%else.mask = extractvalue { i1, i64 } %else, 1
 define amdgpu_ps void @test_else(i64 inreg %mask) {
 entry:
   %else = call { i1, i64 } @llvm.amdgcn.else.i64.i64(i64 %mask)
   %else.bool = extractvalue { i1, i64 } %else, 0
   %else.mask = extractvalue { i1, i64 } %else, 1
   %else.bool.ext = zext i1 %else.bool to i32
   store volatile i32 %else.bool.ext, i32 addrspace(1)* undef
   store volatile i64 %else.mask, i64 addrspace(1)* undef
   ret void
 }
 
 ; This case is probably always broken
-; CHECK: Printing analysis 'Legacy Divergence Analysis' for function 'test_else_divergent_mask':
+; CHECK: Divergence Analysis' for function 'test_else_divergent_mask':
 ; CHECK: DIVERGENT: %if = call { i1, i64 } @llvm.amdgcn.else.i64.i64(i64 %mask)
 ; CHECK-NEXT: DIVERGENT: %if.bool = extractvalue { i1, i64 } %if, 0
 ; CHECK-NOT: DIVERGENT
 ; CHECK: DIVERGENT: %if.bool.ext = zext i1 %if.bool to i32
 define void @test_else_divergent_mask(i64 %mask) {
 entry:
   %if = call { i1, i64 } @llvm.amdgcn.else.i64.i64(i64 %mask)
   %if.bool = extractvalue { i1, i64 } %if, 0

llvm/test/Analysis/DivergenceAnalysis/AMDGPU/hidden_diverge.ll

-; RUN: opt -mtriple amdgcn-unknown-amdhsa -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt -mtriple amdgcn-unknown-amdhsa -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 
 define amdgpu_kernel void @hidden_diverge(i32 %n, i32 %a, i32 %b) #0 {
-; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'hidden_diverge'
+; CHECK-LABEL: 'Divergence Analysis' for function 'hidden_diverge'
 entry:
   %tid = call i32 @llvm.amdgcn.workitem.id.x()
   %cond.var = icmp slt i32 %tid, 0
   br i1 %cond.var, label %B, label %C ; divergent
 ; CHECK: DIVERGENT: br i1 %cond.var,
 B:
   %cond.uni = icmp slt i32 %n, 0
   br i1 %cond.uni, label %C, label %merge ; uniform
 ; CHECK-NOT: DIVERGENT: br i1 %cond.uni,
 C:
   %phi.var.hidden = phi i32 [ 1, %entry ], [ 2, %B  ]
 ; CHECK: DIVERGENT: %phi.var.hidden = phi i32
   br label %merge
 merge:
   %phi.ipd = phi i32 [ %a, %B ], [ %b, %C ]
 ; CHECK: DIVERGENT: %phi.ipd = phi i32
   ret void
 }
 
 define amdgpu_kernel void @hidden_loop_ipd(i32 %n, i32 %a, i32 %b) #0 {
-; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'hidden_loop_ipd'
+; CHECK-LABEL: 'Divergence Analysis' for function 'hidden_loop_ipd'
 entry:
   %tid = call i32 @llvm.amdgcn.workitem.id.x()
   %cond.var = icmp slt i32 %tid, 0
 ; CHECK: DIVERGENT: %cond.var = icmp
   %cond.uni = icmp slt i32 %n, 0
 ; CHECK-NOT: DIVERGENT: %cond.uni = icmp
   br label %for.header
 for.header:

llvm/test/Analysis/DivergenceAnalysis/AMDGPU/hidden_loopdiverge.ll

-; RUN: opt -mtriple amdgcn-unknown-amdhsa -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt -mtriple amdgcn-unknown-amdhsa -enable-new-pm=0 -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt -mtriple amdgcn-unknown-amdhsa -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 
 ; divergent loop (H<header><exiting to X>, B<exiting to Y>)
 ; the divergent join point in %exit is obscured by uniform control joining in %X
 define amdgpu_kernel void @hidden_loop_diverge(i32 %n, i32 %a, i32 %b) #0 {
-; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'hidden_loop_diverge':
+; CHECK-LABEL: Divergence Analysis' for function 'hidden_loop_diverge'
 ; CHECK-NOT: DIVERGENT: %uni.
 ; CHECK-NOT: DIVERGENT: br i1 %uni.
 
 entry:
   %tid = call i32 @llvm.amdgcn.workitem.id.x()
   %uni.cond = icmp slt i32 %a, 0
   br i1 %uni.cond, label %X, label %H  ; uniform
 
exit:
   %div.merge.exit = phi i32 [ %a, %X ], [ %b, %Y ]
   ret void
 ; CHECK: DIVERGENT: %div.merge.exit =
 }
 
 ; divergent loop (H<header><exiting to X>, B<exiting to Y>)
 ; the phi nodes in X and Y don't actually receive divergent values
 define amdgpu_kernel void @unobserved_loop_diverge(i32 %n, i32 %a, i32 %b) #0 {
-; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'unobserved_loop_diverge':
+; CHECK-LABEL: Divergence Analysis' for function 'unobserved_loop_diverge':
 ; CHECK-NOT: DIVERGENT: %uni.
 ; CHECK-NOT: DIVERGENT: br i1 %uni.
 
 entry:
   %tid = call i32 @llvm.amdgcn.workitem.id.x()
   %uni.cond = icmp slt i32 %a, 0
   br i1 %uni.cond, label %X, label %H  ; uniform
 
exit:
   ret void
 ; CHECK: DIVERGENT: %div.merge.exit =
 }
 
 ; divergent loop (G<header>, L<exiting to D>) inside divergent loop (H<header>, B<exiting to X>, C<exiting to Y>, D, G, L)
 ; the inner loop has no exit to top level.
 ; the outer loop becomes divergent as its exiting branch in C is control-dependent on the inner loop's divergent loop exit in D.
 define amdgpu_kernel void @hidden_nestedloop_diverge(i32 %n, i32 %a, i32 %b) #0 {
-; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'hidden_nestedloop_diverge':
+; CHECK-LABEL: Divergence Analysis' for function 'hidden_nestedloop_diverge':
 ; CHECK-NOT: DIVERGENT: %uni.
 ; CHECK-NOT: DIVERGENT: br i1 %uni.
 
 entry:
   %tid = call i32 @llvm.amdgcn.workitem.id.x()
   %uni.cond = icmp slt i32 %a, 0
   %div.exitx = icmp slt i32 %tid, 0
   br i1 %uni.cond, label %X, label %H
exit:
   ret void
 ; CHECK: DIVERGENT: %div.merge.exit =
 }
 
 ; divergent loop (G<header>, L<exiting to X>) in divergent loop (H<header>, B<exiting to C>, C, G, L)
 ; the outer loop has no immediately divergent exiting edge.
 ; the inner exiting edge is exiting to top-level through the outer loop causing both to become divergent.
 define amdgpu_kernel void @hidden_doublebreak_diverge(i32 %n, i32 %a, i32 %b) #0 {
-; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'hidden_doublebreak_diverge':
+; CHECK-LABEL: Divergence Analysis' for function 'hidden_doublebreak_diverge':
 ; CHECK-NOT: DIVERGENT: %uni.
 ; CHECK-NOT: DIVERGENT: br i1 %uni.
 
 entry:
   %tid = call i32 @llvm.amdgcn.workitem.id.x()
   %uni.cond = icmp slt i32 %a, 0
   %div.exitx = icmp slt i32 %tid, 0
   br i1 %uni.cond, label %X, label %H

 Y:
   %div.merge.y = phi i32 [ 42, %X ], [ %b, %B ]
   ret void
 ; CHECK: DIVERGENT: %div.merge.y =
 }
 
 ; divergent loop (G<header>, L<exiting to D>) contained inside a uniform loop (H<header>, B, G, L , D<exiting to x>)
 define amdgpu_kernel void @hidden_containedloop_diverge(i32 %n, i32 %a, i32 %b) #0 {
-; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'hidden_containedloop_diverge':
+; CHECK-LABEL: Divergence Analysis' for function 'hidden_containedloop_diverge':
 ; CHECK-NOT: DIVERGENT: %uni.
 ; CHECK-NOT: DIVERGENT: br i1 %uni.
 
 entry:
   %tid = call i32 @llvm.amdgcn.workitem.id.x()
   %uni.cond = icmp slt i32 %a, 0
   %div.exitx = icmp slt i32 %tid, 0
   br i1 %uni.cond, label %X, label %H

llvm/test/Analysis/DivergenceAnalysis/AMDGPU/inline-asm.ll

-; RUN: opt -mtriple=amdgcn-unknown-amdhsa -mcpu=tahiti -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
-; RUN: opt -mtriple=amdgcn-unknown-amdhsa -mcpu=gfx908 -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt -mtriple=amdgcn-unknown-amdhsa -mcpu=tahiti -enable-new-pm=0 -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt -mtriple=amdgcn-unknown-amdhsa -mcpu=gfx908 -enable-new-pm=0 -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt -mtriple amdgcn-unknown-amdhsa -mcpu=tahiti -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
+; RUN: opt -mtriple amdgcn-unknown-amdhsa -mcpu=gfx908 -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 ; Make sure nothing crashes on targets with or without AGPRs
 
-; CHECK: Printing analysis 'Legacy Divergence Analysis' for function 'inline_asm_1_sgpr_virtreg_output':
+; CHECK: Divergence Analysis' for function 'inline_asm_1_sgpr_virtreg_output':
 ; CHECK-NOT: DIVERGENT
 define i32 @inline_asm_1_sgpr_virtreg_output() {
   %sgpr = call i32 asm "s_mov_b32 $0, 0", "=s"()
   ret i32 %sgpr
 }
 
-; CHECK: Printing analysis 'Legacy Divergence Analysis' for function 'inline_asm_1_sgpr_physreg_output':
+; CHECK: Divergence Analysis' for function 'inline_asm_1_sgpr_physreg_output':
 ; CHECK-NOT: DIVERGENT
 define i32 @inline_asm_1_sgpr_physreg_output() {
   %sgpr = call i32 asm "s_mov_b32 s0, 0", "={s0}"()
   ret i32 %sgpr
 }
 
-; CHECK: Printing analysis 'Legacy Divergence Analysis' for function 'inline_asm_1_vgpr_virtreg_output':
+; CHECK: Divergence Analysis' for function 'inline_asm_1_vgpr_virtreg_output':
 ; CHECK: DIVERGENT: %vgpr = call i32 asm "v_mov_b32 $0, 0", "=v"()
 define i32 @inline_asm_1_vgpr_virtreg_output() {
   %vgpr = call i32 asm "v_mov_b32 $0, 0", "=v"()
   ret i32 %vgpr
 }
 
-; CHECK: Printing analysis 'Legacy Divergence Analysis' for function 'inline_asm_1_vgpr_physreg_output':
+; CHECK: Divergence Analysis' for function 'inline_asm_1_vgpr_physreg_output':
 ; CHECK: DIVERGENT: %vgpr = call i32 asm "v_mov_b32 v0, 0", "={v0}"()
 define i32 @inline_asm_1_vgpr_physreg_output() {
   %vgpr = call i32 asm "v_mov_b32 v0, 0", "={v0}"()
   ret i32 %vgpr
 }
 
-; CHECK: Printing analysis 'Legacy Divergence Analysis' for function 'inline_asm_1_agpr_virtreg_output':
+; CHECK: Divergence Analysis' for function 'inline_asm_1_agpr_virtreg_output':
 ; CHECK: DIVERGENT: %vgpr = call i32 asm "; def $0", "=a"()
 define i32 @inline_asm_1_agpr_virtreg_output() {
   %vgpr = call i32 asm "; def $0", "=a"()
   ret i32 %vgpr
 }
 
-; CHECK: Printing analysis 'Legacy Divergence Analysis' for function 'inline_asm_1_agpr_physreg_output':
+; CHECK: Divergence Analysis' for function 'inline_asm_1_agpr_physreg_output':
 ; CHECK: DIVERGENT: %vgpr = call i32 asm "; def a0", "={a0}"()
 define i32 @inline_asm_1_agpr_physreg_output() {
   %vgpr = call i32 asm "; def a0", "={a0}"()
   ret i32 %vgpr
 }
 
-; CHECK: Printing analysis 'Legacy Divergence Analysis' for function 'inline_asm_2_sgpr_virtreg_output':
+; CHECK: Divergence Analysis' for function 'inline_asm_2_sgpr_virtreg_output':
 ; CHECK-NOT: DIVERGENT
 define void @inline_asm_2_sgpr_virtreg_output() {
   %asm = call { i32, i32 } asm "; def $0, $1", "=s,=s"()
   %sgpr0 = extractvalue { i32, i32 } %asm, 0
   %sgpr1 = extractvalue { i32, i32 } %asm, 1
   store i32 %sgpr0, i32 addrspace(1)* undef
   store i32 %sgpr1, i32 addrspace(1)* undef
   ret void
 }
 
 ; One output is SGPR, one is VGPR. Infer divergent for the aggregate, but uniform on the SGPR extract
-; CHECK: Printing analysis 'Legacy Divergence Analysis' for function 'inline_asm_sgpr_vgpr_virtreg_output':
+; CHECK: Divergence Analysis' for function 'inline_asm_sgpr_vgpr_virtreg_output':
 ; CHECK: DIVERGENT:       %asm = call { i32, i32 } asm "; def $0, $1", "=s,=v"()
 ; CHECK-NEXT: {{^[ \t]+}}%sgpr = extractvalue { i32, i32 } %asm, 0
 ; CHECK-NEXT: DIVERGENT:       %vgpr = extractvalue { i32, i32 } %asm, 1
 define void @inline_asm_sgpr_vgpr_virtreg_output() {
   %asm = call { i32, i32 } asm "; def $0, $1", "=s,=v"()
   %sgpr = extractvalue { i32, i32 } %asm, 0
   %vgpr = extractvalue { i32, i32 } %asm, 1
   store i32 %sgpr, i32 addrspace(1)* undef
   store i32 %vgpr, i32 addrspace(1)* undef
   ret void
 }
 
-; CHECK: Printing analysis 'Legacy Divergence Analysis' for function 'inline_asm_vgpr_sgpr_virtreg_output':
+; CHECK: Divergence Analysis' for function 'inline_asm_vgpr_sgpr_virtreg_output':
 ; CHECK: DIVERGENT:       %asm = call { i32, i32 } asm "; def $0, $1", "=v,=s"()
 ; CHECK-NEXT: DIVERGENT:       %vgpr = extractvalue { i32, i32 } %asm, 0
 ; CHECK-NEXT: {{^[ \t]+}}%sgpr = extractvalue { i32, i32 } %asm, 1
 define void @inline_asm_vgpr_sgpr_virtreg_output() {
   %asm = call { i32, i32 } asm "; def $0, $1", "=v,=s"()
   %vgpr = extractvalue { i32, i32 } %asm, 0
   %sgpr = extractvalue { i32, i32 } %asm, 1
   store i32 %vgpr, i32 addrspace(1)* undef
   store i32 %sgpr, i32 addrspace(1)* undef
   ret void
 }
 
 ; Have an extra output constraint
-; CHECK: Printing analysis 'Legacy Divergence Analysis' for function 'multi_sgpr_inline_asm_output_input_constraint':
+; CHECK: Divergence Analysis' for function 'multi_sgpr_inline_asm_output_input_constraint':
 ; CHECK-NOT: DIVERGENT
 define void @multi_sgpr_inline_asm_output_input_constraint() {
   %asm = call { i32, i32 } asm "; def $0, $1", "=s,=s,s"(i32 1234)
   %sgpr0 = extractvalue { i32, i32 } %asm, 0
   %sgpr1 = extractvalue { i32, i32 } %asm, 1
   store i32 %sgpr0, i32 addrspace(1)* undef
   store i32 %sgpr1, i32 addrspace(1)* undef
   ret void
 }
 
-; CHECK: Printing analysis 'Legacy Divergence Analysis' for function 'inline_asm_vgpr_sgpr_virtreg_output_input_constraint':
+; CHECK: Divergence Analysis' for function 'inline_asm_vgpr_sgpr_virtreg_output_input_constraint':
 ; CHECK: DIVERGENT:       %asm = call { i32, i32 } asm "; def $0, $1", "=v,=s,v"(i32 1234)
 ; CHECK-NEXT: DIVERGENT:       %vgpr = extractvalue { i32, i32 } %asm, 0
 ; CHECK-NEXT: {{^[ \t]+}}%sgpr = extractvalue { i32, i32 } %asm, 1
 define void @inline_asm_vgpr_sgpr_virtreg_output_input_constraint() {
   %asm = call { i32, i32 } asm "; def $0, $1", "=v,=s,v"(i32 1234)
   %vgpr = extractvalue { i32, i32 } %asm, 0
   %sgpr = extractvalue { i32, i32 } %asm, 1
   store i32 %vgpr, i32 addrspace(1)* undef
   store i32 %sgpr, i32 addrspace(1)* undef
   ret void
 }

llvm/test/Analysis/DivergenceAnalysis/AMDGPU/interp_f16.ll

-; RUN: opt -mtriple=amdgcn-- -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt -mtriple=amdgcn-- -enable-new-pm=0 -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt -mtriple amdgcn-- -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 
 ; CHECK: for function 'interp_p1_f16'
 ; CHECK: DIVERGENT:       %p1 = call float @llvm.amdgcn.interp.p1.f16
 define amdgpu_ps float @interp_p1_f16(float inreg %i, float inreg %j, i32 inreg %m0) #0 {
 main_body:
   %p1 = call float @llvm.amdgcn.interp.p1.f16(float %i, i32 1, i32 2, i1 0, i32 %m0)
   ret float %p1
 }

llvm/test/Analysis/DivergenceAnalysis/AMDGPU/intrinsics.ll

-; RUN: opt -mtriple=amdgcn-- -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt -mtriple=amdgcn-- -enable-new-pm=0 -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt -mtriple amdgcn-- -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 
 ; CHECK: DIVERGENT: %swizzle = call i32 @llvm.amdgcn.ds.swizzle(i32 %src, i32 100) #0
 define amdgpu_kernel void @ds_swizzle(i32 addrspace(1)* %out, i32 %src) #0 {
   %swizzle = call i32 @llvm.amdgcn.ds.swizzle(i32 %src, i32 100) #0
   store i32 %swizzle, i32 addrspace(1)* %out, align 4
   ret void
 }
 

llvm/test/Analysis/DivergenceAnalysis/AMDGPU/irreducible.ll

-; RUN: opt -mtriple=amdgcn-- -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt %s -mtriple amdgcn-- -enable-new-pm=0 -analyze -divergence -use-gpu-divergence-analysis | FileCheck %s
+; RUN: opt %s -mtriple amdgcn-- -passes='print<divergence>' -disable-output 2>&1 | FileCheck %s
+
+; NOTE: The new pass manager does not fall back on legacy divergence
+; analysis even when the function contains an irreducible loop. The
+; (new) divergence analysis conservatively reports all values as
+; divergent. This test does not check for this conservative
+; behaviour. Instead, it only checks for the values that are known to
+; be divergent according to the legacy analysis.
+
+; RUN: opt -mtriple amdgcn-- -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 
 ; This test contains an unstructured loop.
 ;           +-------------- entry ----------------+
 ;           |                                     |
 ;           V                                     V
 ; i1 = phi(0, i3)                            i2 = phi(0, i3)
 ;     j1 = i1 + 1 ---> i3 = phi(j1, j2) <--- j2 = i2 + 2
 ;           ^                 |                   ^
 ;           |                 V                   |
 ;           +-------- switch (tid / i3) ----------+
 ;                             |
 ;                             V
 ;                        if (i3 == 5) // divergent
 ; because sync dependent on (tid / i3).
 define i32 @unstructured_loop(i1 %entry_cond) {
-; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'unstructured_loop'
+; CHECK-LABEL: Divergence Analysis' for function 'unstructured_loop'
 entry:
   %tid = call i32 @llvm.amdgcn.workitem.id.x()
   br i1 %entry_cond, label %loop_entry_1, label %loop_entry_2
 loop_entry_1:
   %i1 = phi i32 [ 0, %entry ], [ %i3, %loop_latch ]
   %j1 = add i32 %i1, 1
   br label %loop_body
 loop_entry_2:

llvm/test/Analysis/DivergenceAnalysis/AMDGPU/join-at-loop-exit.ll

-; RUN: opt -mtriple amdgcn-unknown-amdhsa -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt -mtriple amdgcn-unknown-amdhsa -enable-new-pm=0 -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt -mtriple amdgcn-unknown-amdhsa -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 
 ; CHECK: bb3:
 ; CHECK: DIVERGENT:       %Guard.bb4 = phi i1 [ true, %bb1 ], [ false, %bb2 ]
 ; CHECK: DIVERGENT:       br i1 %Guard.bb4, label %bb4, label %bb5
 
 ; Function Attrs: nounwind readnone speculatable
 declare i32 @llvm.amdgcn.workitem.id.x() #0
 

llvm/test/Analysis/DivergenceAnalysis/AMDGPU/kernel-args.ll

-; RUN: opt %s -mtriple amdgcn-- -analyze -divergence -use-gpu-divergence-analysis | FileCheck %s
+; RUN: opt %s -mtriple amdgcn-- -enable-new-pm=0 -analyze -divergence -use-gpu-divergence-analysis | FileCheck %s
+; RUN: opt -mtriple amdgcn-- -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 
-; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'test_amdgpu_ps':
+; CHECK-LABEL: Divergence Analysis' for function 'test_amdgpu_ps':
 ; CHECK: DIVERGENT:  [4 x <16 x i8>] addrspace(4)* %arg0
 ; CHECK-NOT: DIVERGENT
 ; CHECK: DIVERGENT:  <2 x i32> %arg3
 ; CHECK: DIVERGENT:  <3 x i32> %arg4
 ; CHECK: DIVERGENT:  float %arg5
 ; CHECK: DIVERGENT:  i32 %arg6
 
 define amdgpu_ps void @test_amdgpu_ps([4 x <16 x i8>] addrspace(4)* byref([4 x <16 x i8>]) %arg0, float inreg %arg1, i32 inreg %arg2, <2 x i32> %arg3, <3 x i32> %arg4, float %arg5, i32 %arg6) #0 {
   ret void
 }
 
-; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'test_amdgpu_kernel':
+; CHECK-LABEL: Divergence Analysis' for function 'test_amdgpu_kernel':
 ; CHECK-NOT: %arg0
 ; CHECK-NOT: %arg1
 ; CHECK-NOT: %arg2
 ; CHECK-NOT: %arg3
 ; CHECK-NOT: %arg4
 ; CHECK-NOT: %arg5
 ; CHECK-NOT: %arg6
 define amdgpu_kernel void @test_amdgpu_kernel([4 x <16 x i8>] addrspace(4)* byref([4 x <16 x i8>]) %arg0, float inreg %arg1, i32 inreg %arg2, <2 x i32> %arg3, <3 x i32> %arg4, float %arg5, i32 %arg6) #0 {
   ret void
 }
 
-; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'test_c':
+; CHECK-LABEL: Divergence Analysis' for function 'test_c':
 ; CHECK: DIVERGENT:
 ; CHECK: DIVERGENT:
 ; CHECK: DIVERGENT:
 ; CHECK: DIVERGENT:
 ; CHECK: DIVERGENT:
 ; CHECK: DIVERGENT:
 ; CHECK: DIVERGENT:
 define void @test_c([4 x <16 x i8>] addrspace(5)* byval([4 x <16 x i8>]) %arg0, float inreg %arg1, i32 inreg %arg2, <2 x i32> %arg3, <3 x i32> %arg4, float %arg5, i32 %arg6) #0 {
   ret void
 }
 
 attributes #0 = { nounwind }

llvm/test/Analysis/DivergenceAnalysis/AMDGPU/llvm.amdgcn.buffer.atomic.ll

-;RUN: opt -mtriple=amdgcn-mesa-mesa3d -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+;RUN: opt -mtriple=amdgcn-mesa-mesa3d -enable-new-pm=0 -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt -mtriple amdgcn-mesa-mesa3d -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 
 ;CHECK: DIVERGENT: %orig = call i32 @llvm.amdgcn.buffer.atomic.swap.i32(
 define float @buffer_atomic_swap(<4 x i32> inreg %rsrc, i32 inreg %data) #0 {
 main_body:
   %orig = call i32 @llvm.amdgcn.buffer.atomic.swap.i32(i32 %data, <4 x i32> %rsrc, i32 0, i32 0, i1 0)
   %r = bitcast i32 %orig to float
   ret float %r
 }

llvm/test/Analysis/DivergenceAnalysis/AMDGPU/llvm.amdgcn.image.atomic.ll

-;RUN: opt -mtriple=amdgcn-mesa-mesa3d -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+;RUN: opt -mtriple=amdgcn-mesa-mesa3d -enable-new-pm=0 -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt -mtriple amdgcn-mesa-mesa3d -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 
 ;CHECK: DIVERGENT: %orig = call i32 @llvm.amdgcn.image.atomic.swap.1d.i32.i32(
 define float @image_atomic_swap(<8 x i32> inreg %rsrc, i32 inreg %addr, i32 inreg %data) #0 {
 main_body:
   %orig = call i32 @llvm.amdgcn.image.atomic.swap.1d.i32.i32(i32 %data, i32 %addr, <8 x i32> %rsrc, i32 0, i32 0)
   %r = bitcast i32 %orig to float
   ret float %r
 }

llvm/test/Analysis/DivergenceAnalysis/AMDGPU/no-return-blocks.ll

-; RUN: opt %s -mtriple amdgcn-- -analyze -divergence -use-gpu-divergence-analysis | FileCheck %s
+; RUN: opt %s -mtriple amdgcn-- -enable-new-pm=0 -analyze -divergence -use-gpu-divergence-analysis | FileCheck %s
+; RUN: opt -mtriple amdgcn-- -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 
 ; CHECK: DIVERGENT:  %tmp5 = getelementptr inbounds float, float addrspace(1)* %arg, i64 %tmp2
 ; CHECK: DIVERGENT:  %tmp10 = load volatile float, float addrspace(1)* %tmp5, align 4
 ; CHECK: DIVERGENT:  %tmp11 = load volatile float, float addrspace(1)* %tmp5, align 4
 
 ; The post dominator tree does not have a root node in this case
 define amdgpu_kernel void @no_return_blocks(float addrspace(1)* noalias nocapture readonly %arg, float addrspace(1)* noalias nocapture readonly %arg1) #0 {
 bb0:

llvm/test/Analysis/DivergenceAnalysis/AMDGPU/phi-undef.ll

-; RUN: opt -mtriple=amdgcn-- -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt -mtriple=amdgcn-- -enable-new-pm=0 -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt -mtriple amdgcn-- -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 
 ; CHECK-LABEL: 'test1':
 ; CHECK-NEXT: DIVERGENT: i32 %bound
 ; CHECK: {{^  *}}%counter =
 ; CHECK-NEXT: DIVERGENT: %break = icmp sge i32 %counter, %bound
 ; CHECK-NEXT: DIVERGENT: br i1 %break, label %footer, label %body
 ; CHECK: {{^  *}}%counter.next =
 ; CHECK: {{^  *}}%counter.footer =

llvm/test/Analysis/DivergenceAnalysis/AMDGPU/propagate-loop-live-out.ll

-; RUN: opt -mtriple amdgcn-unknown-amdhsa -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt -mtriple amdgcn-unknown-amdhsa -enable-new-pm=0 -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt -mtriple amdgcn-unknown-amdhsa -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 
 ; CHECK: bb6:
 ; CHECK: DIVERGENT:       %.126.i355.i = phi i1 [ false, %bb5 ], [ true, %bb4 ]
 ; CHECK: DIVERGENT:       br i1 %.126.i355.i, label %bb7, label %bb8
 
 ; Function Attrs: nounwind readnone speculatable
 declare i32 @llvm.amdgcn.workitem.id.x() #0
 

llvm/test/Analysis/DivergenceAnalysis/AMDGPU/temporal_diverge.ll

-; RUN: opt -mtriple amdgcn-unknown-amdhsa -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt -mtriple amdgcn-unknown-amdhsa -enable-new-pm=0 -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt -mtriple amdgcn-unknown-amdhsa -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 
 ; temporal-divergent use of value carried by divergent loop
 define amdgpu_kernel void @temporal_diverge(i32 %n, i32 %a, i32 %b) #0 {
-; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'temporal_diverge':
+; CHECK-LABEL: Divergence Analysis' for function 'temporal_diverge':
 ; CHECK-NOT: DIVERGENT: %uni.
 ; CHECK-NOT: DIVERGENT: br i1 %uni.
 
 entry:
   %tid = call i32 @llvm.amdgcn.workitem.id.x()
   %uni.cond = icmp slt i32 %a, 0
   br label %H
 
 H:
   %uni.merge.h = phi i32 [ 0, %entry ], [ %uni.inc, %H ]
   %uni.inc = add i32 %uni.merge.h, 1
   %div.exitx = icmp slt i32 %tid, 0
   br i1 %div.exitx, label %X, label %H ; divergent branch
 ; CHECK: DIVERGENT: %div.exitx =  
 ; CHECK: DIVERGENT: br i1 %div.exitx, 
 
 X:
   %div.user = add i32 %uni.inc, 5
   ret void
 }
 
 ; temporal-divergent use of value carried by divergent loop inside a top-level loop
 define amdgpu_kernel void @temporal_diverge_inloop(i32 %n, i32 %a, i32 %b) #0 {
-; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'temporal_diverge_inloop':
+; CHECK-LABEL: Divergence Analysis' for function 'temporal_diverge_inloop':
 ; CHECK-NOT: DIVERGENT: %uni.
 ; CHECK-NOT: DIVERGENT: br i1 %uni.
 
 entry:
   %tid = call i32 @llvm.amdgcn.workitem.id.x()
   %uni.cond = icmp slt i32 %a, 0
   br label %G
 

 Y:
   %div.alsouser = add i32 %uni.inc, 5
   ret void
 }
 
 
 ; temporal-uniform use of a valud, definition and users are carried by a surrounding divergent loop
 define amdgpu_kernel void @temporal_uniform_indivloop(i32 %n, i32 %a, i32 %b) #0 {
-; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'temporal_uniform_indivloop':
+; CHECK-LABEL: Divergence Analysis' for function 'temporal_uniform_indivloop':
 ; CHECK-NOT: DIVERGENT: %uni.
 ; CHECK-NOT: DIVERGENT: br i1 %uni.
 
 entry:
   %tid = call i32 @llvm.amdgcn.workitem.id.x()
   %uni.cond = icmp slt i32 %a, 0
   br label %G
 

 Y:
   %div.alsouser = add i32 %uni.inc, 5
   ret void
 }
 
 
 ; temporal-divergent use of value carried by divergent loop, user is inside sibling loop
 define amdgpu_kernel void @temporal_diverge_loopuser(i32 %n, i32 %a, i32 %b) #0 {
-; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'temporal_diverge_loopuser':
+; CHECK-LABEL: Divergence Analysis' for function 'temporal_diverge_loopuser':
 ; CHECK-NOT: DIVERGENT: %uni.
 ; CHECK-NOT: DIVERGENT: br i1 %uni.
 
 entry:
   %tid = call i32 @llvm.amdgcn.workitem.id.x()
   %uni.cond = icmp slt i32 %a, 0
   br label %H
 
G:
   br i1 %uni.cond, label %G, label %Y
 
 Y:
   ret void
 }
 
 ; temporal-divergent use of value carried by divergent loop, user is inside sibling loop, defs and use are carried by a uniform loop
 define amdgpu_kernel void @temporal_diverge_loopuser_nested(i32 %n, i32 %a, i32 %b) #0 {
-; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'temporal_diverge_loopuser_nested':
+; CHECK-LABEL: Divergence Analysis' for function 'temporal_diverge_loopuser_nested':
 ; CHECK-NOT: DIVERGENT: %uni.
 ; CHECK-NOT: DIVERGENT: br i1 %uni.
 
 entry:
   %tid = call i32 @llvm.amdgcn.workitem.id.x()
   %uni.cond = icmp slt i32 %a, 0
   br label %H
 

llvm/test/Analysis/DivergenceAnalysis/AMDGPU/trivial-join-at-loop-exit.ll

-; RUN: opt -mtriple amdgcn-unknown-amdhsa -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt -mtriple amdgcn-unknown-amdhsa -enable-new-pm=0 -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt -mtriple amdgcn-unknown-amdhsa -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 
 ; CHECK: bb2:
 ; CHECK-NOT: DIVERGENT:       %Guard.bb2 = phi i1 [ true, %bb1 ], [ false, %bb0 ]
 
 ; Function Attrs: nounwind readnone speculatable
 declare i32 @llvm.amdgcn.workitem.id.x() #0
 
 define protected amdgpu_kernel void @test2(i1 %uni) {

llvm/test/Analysis/DivergenceAnalysis/AMDGPU/unreachable-loop-block.ll

-; RUN: opt %s -mtriple amdgcn-- -analyze -divergence -use-gpu-divergence-analysis | FileCheck %s
+; RUN: opt %s -mtriple amdgcn-- -enable-new-pm=0 -analyze -divergence -use-gpu-divergence-analysis | FileCheck %s
+; RUN: opt -mtriple amdgcn-- -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 
 ; CHECK: DIVERGENT:  %tmp = cmpxchg volatile
 define amdgpu_kernel void @unreachable_loop(i32 %tidx) #0 {
 entry:
   unreachable
 
 unreachable_loop:                                        ; preds = %do.body.i, %if.then11
   %tmp = cmpxchg volatile i32 addrspace(1)* null, i32 0, i32 0 seq_cst seq_cst
   %cmp.i = extractvalue { i32, i1 } %tmp, 1
   br i1 %cmp.i, label %unreachable_loop, label %end
 
 end:                                      ; preds = %do.body.i51, %atomicAdd_g_f.exit
   unreachable
 }
 
 attributes #0 = { norecurse nounwind }

llvm/test/Analysis/DivergenceAnalysis/AMDGPU/workitem-intrinsics.ll

-; RUN: opt  -mtriple amdgcn-unknown-amdhsa -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt  -mtriple amdgcn-unknown-amdhsa -enable-new-pm=0 -analyze -divergence -use-gpu-divergence-analysis %s | FileCheck %s
+; RUN: opt -mtriple amdgcn-unknown-amdhsa -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 
 declare i32 @llvm.amdgcn.workitem.id.x() #0
 declare i32 @llvm.amdgcn.workitem.id.y() #0
 declare i32 @llvm.amdgcn.workitem.id.z() #0
 declare i32 @llvm.amdgcn.mbcnt.lo(i32, i32) #0
 declare i32 @llvm.amdgcn.mbcnt.hi(i32, i32) #0
 
 ; CHECK: DIVERGENT:  %id.x = call i32 @llvm.amdgcn.workitem.id.x()

llvm/test/Analysis/DivergenceAnalysis/NVPTX/daorder.ll

-; RUN: opt %s -analyze -divergence -use-gpu-divergence-analysis | FileCheck %s
+; RUN: opt %s -enable-new-pm=0 -analyze -divergence -use-gpu-divergence-analysis | FileCheck %s
+; RUN: opt %s -passes='print<divergence>' -disable-output 2>&1 | FileCheck %s
 
 target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"
 target triple = "nvptx64-nvidia-cuda"
 
 define i32 @daorder(i32 %n) {
-; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'daorder'
+; CHECK-LABEL: Divergence Analysis' for function 'daorder'
 entry:
   %tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
   %cond = icmp slt i32 %tid, 0
   br i1 %cond, label %A, label %B ; divergent
 ; CHECK: DIVERGENT: br i1 %cond,
 A:
   %defAtA = add i32 %n, 1 ; uniform
 ; CHECK-NOT: DIVERGENT: %defAtA = 

llvm/test/Analysis/DivergenceAnalysis/NVPTX/diverge.ll

-; RUN: opt %s -analyze -divergence -use-gpu-divergence-analysis | FileCheck %s
+; RUN: opt %s -enable-new-pm=0 -analyze -divergence -use-gpu-divergence-analysis | FileCheck %s
+; RUN: opt %s -passes='print<divergence>' -disable-output 2>&1 | FileCheck %s
 
 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 'Legacy Divergence Analysis' for function 'no_diverge'
+; CHECK-LABEL: Divergence Analysis' for function 'no_diverge'
 entry:
   %tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
   %cond = icmp slt i32 %n, 0
   br i1 %cond, label %then, label %else ; uniform
 ; CHECK-NOT: DIVERGENT: br i1 %cond,
 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;               // c is divergent: sync dependent
 define i32 @sync(i32 %a, i32 %b) {
-; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'sync'
+; CHECK-LABEL: Divergence Analysis' for function 'sync'
 bb1:
   %tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.y()
   %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
 ; CHECK: DIVERGENT: %c =
   ret i32 %c
 }
 
 ; c = 0;
 ; if (threadIdx.x >= 5) {  // divergent
 ;   c = (n < 0 ? a : b);  // c here is uniform because n is uniform
 ; }
 ; // c here is divergent because it is sync dependent on threadIdx.x >= 5
 ; return c;
 define i32 @mixed(i32 %n, i32 %a, i32 %b) {
-; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'mixed'
+; CHECK-LABEL: Divergence Analysis' for function 'mixed'
 bb1:
   %tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.z()
   %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
 bb3:
   br label %bb5
 bb4:
   br label %bb5
 bb5:
   %c = phi i32 [ %a, %bb3 ], [ %b, %bb4 ]
 ; CHECK-NOT: DIVERGENT: %c =
   br label %bb6
 bb6:
   %c2 = phi i32 [ 0, %bb1], [ %c, %bb5 ]
 ; CHECK: DIVERGENT: %c2 =
   ret i32 %c2
 }
 
 ; We conservatively treats all parameters of a __device__ function as divergent.
 define i32 @device(i32 %n, i32 %a, i32 %b) {
-; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'device'
+; CHECK-LABEL: Divergence Analysis' for function 'device'
 ; CHECK: DIVERGENT: i32 %n
 ; CHECK: DIVERGENT: i32 %a
 ; CHECK: DIVERGENT: i32 %b
 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
 }
 
 ; int i = 0;
 ; do {
 ;   i++;                  // i here is uniform
 ; } while (i < laneid);
 ; return i == 10 ? 0 : 1; // i here is divergent
 ;
 ; The i defined in the loop is used outside.
 define i32 @loop() {
-; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'loop'
+; CHECK-LABEL: Divergence Analysis' for function 'loop'
 entry:
   %laneid = call i32 @llvm.nvvm.read.ptx.sreg.laneid()
   br label %loop
 loop:
   %i = phi i32 [ 0, %entry ], [ %i1, %loop ]
 ; CHECK-NOT: DIVERGENT: %i =
   %i1 = add i32 %i, 1
   %exit_cond = icmp sge i32 %i1, %laneid
   br i1 %exit_cond, label %loop_exit, label %loop
 loop_exit:
   %cond = icmp eq i32 %i, 10
   br i1 %cond, label %then, label %else
 ; CHECK: DIVERGENT: br i1 %cond,
 then:
   ret i32 0
 else:
   ret i32 1
 }
 
 ; Same as @loop, but the loop is in the LCSSA form.
 define i32 @lcssa() {
-; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'lcssa'
+; CHECK-LABEL: Divergence Analysis' for function 'lcssa'
 entry:
   %tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
   br label %loop
 loop:
   %i = phi i32 [ 0, %entry ], [ %i1, %loop ]
 ; CHECK-NOT: DIVERGENT: %i =
   %i1 = add i32 %i, 1
   %exit_cond = icmp sge i32 %i1, %tid

llvm/test/Analysis/DivergenceAnalysis/NVPTX/hidden_diverge.ll

-; RUN: opt %s -analyze -divergence -use-gpu-divergence-analysis | FileCheck %s
+; RUN: opt %s -enable-new-pm=0 -analyze -divergence -use-gpu-divergence-analysis | FileCheck %s
+; RUN: opt %s -passes='print<divergence>' -disable-output 2>&1 | FileCheck %s
 
 target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"
 target triple = "nvptx64-nvidia-cuda"
 
 define i32 @hidden_diverge(i32 %n, i32 %a, i32 %b) {
-; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'hidden_diverge'
+; CHECK-LABEL: Divergence Analysis' for function 'hidden_diverge'
 entry:
   %tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
   %cond.var = icmp slt i32 %tid, 0
   br i1 %cond.var, label %B, label %C ; divergent
 ; CHECK: DIVERGENT: br i1 %cond.var,
 B:
   %cond.uni = icmp slt i32 %n, 0
   br i1 %cond.uni, label %C, label %merge ; uniform

llvm/test/Analysis/DivergenceAnalysis/NVPTX/irreducible.ll

-; RUN: opt %s -analyze -divergence -use-gpu-divergence-analysis | FileCheck %s
+; RUN: opt %s -enable-new-pm=0 -analyze -divergence -use-gpu-divergence-analysis | FileCheck %s
+; RUN: opt %s -passes='print<divergence>' -disable-output 2>&1 | FileCheck %s
+
+; NOTE: The new pass manager does not fall back on legacy divergence
+; analysis even when the function contains an irreducible loop. The
+; (new) divergence analysis conservatively reports all values as
+; divergent. This test does not check for this conservative
+; behaviour. Instead, it only checks for the values that are known to
+; be divergent according to the legacy analysis.
 
 target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"
 target triple = "nvptx64-nvidia-cuda"
 
 ; This test contains an unstructured loop.
 ;           +-------------- entry ----------------+
 ;           |                                     |
 ;           V                                     V
 ; i1 = phi(0, i3)                            i2 = phi(0, i3)
 ;     j1 = i1 + 1 ---> i3 = phi(j1, j2) <--- j2 = i2 + 2
 ;           ^                 |                   ^
 ;           |                 V                   |
 ;           +-------- switch (tid / i3) ----------+
 ;                             |
 ;                             V
 ;                        if (i3 == 5) // divergent
 ; because sync dependent on (tid / i3).
 define i32 @unstructured_loop(i1 %entry_cond) {
-; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'unstructured_loop'
+; CHECK-LABEL: Divergence Analysis' for function 'unstructured_loop'
 entry:
   %tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
   br i1 %entry_cond, label %loop_entry_1, label %loop_entry_2
 loop_entry_1:
   %i1 = phi i32 [ 0, %entry ], [ %i3, %loop_latch ]
   %j1 = add i32 %i1, 1
   br label %loop_body
 loop_entry_2:

llvm/test/Analysis/LegacyDivergenceAnalysis/AMDGPU/atomics.ll

-; RUN: opt -mtriple=amdgcn-- -amdgpu-use-legacy-divergence-analysis -analyze -divergence %s | FileCheck %s
+; RUN: opt -mtriple=amdgcn-- -amdgpu-use-legacy-divergence-analysis -enable-new-pm=0 -analyze -divergence %s | FileCheck %s
 
 ; CHECK: DIVERGENT: %orig = atomicrmw xchg i32* %ptr, i32 %val seq_cst
 define i32 @test1(i32* %ptr, i32 %val) #0 {
   %orig = atomicrmw xchg i32* %ptr, i32 %val seq_cst
   ret i32 %orig
 }
 
 ; CHECK: DIVERGENT: %orig = cmpxchg i32* %ptr, i32 %cmp, i32 %new seq_cst seq_cst

llvm/test/Analysis/LegacyDivergenceAnalysis/AMDGPU/intrinsics.ll

-; RUN: opt -mtriple=amdgcn-- -analyze -amdgpu-use-legacy-divergence-analysis -divergence %s | FileCheck %s
+; RUN: opt -mtriple=amdgcn-- -enable-new-pm=0 -analyze -amdgpu-use-legacy-divergence-analysis -divergence %s | FileCheck %s
 
 ; CHECK: DIVERGENT: %swizzle = call i32 @llvm.amdgcn.ds.swizzle(i32 %src, i32 100) #0
 define amdgpu_kernel void @ds_swizzle(i32 addrspace(1)* %out, i32 %src) #0 {
   %swizzle = call i32 @llvm.amdgcn.ds.swizzle(i32 %src, i32 100) #0
   store i32 %swizzle, i32 addrspace(1)* %out, align 4
   ret void
 }
 
 declare i32 @llvm.amdgcn.ds.swizzle(i32, i32) #1
 
 attributes #0 = { nounwind convergent }
 attributes #1 = { nounwind readnone convergent }

llvm/test/Analysis/LegacyDivergenceAnalysis/AMDGPU/kernel-args.ll

-; RUN: opt %s -mtriple amdgcn-- -amdgpu-use-legacy-divergence-analysis -analyze -divergence | FileCheck %s
+; RUN: opt %s -mtriple amdgcn-- -amdgpu-use-legacy-divergence-analysis -enable-new-pm=0 -analyze -divergence | FileCheck %s
 
 ; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'test_amdgpu_ps':
 ; CHECK: DIVERGENT:  [4 x <16 x i8>] addrspace(4)* %arg0
 ; CHECK-NOT: DIVERGENT
 ; CHECK: DIVERGENT:  <2 x i32> %arg3
 ; CHECK: DIVERGENT:  <3 x i32> %arg4
 ; CHECK: DIVERGENT:  float %arg5
 ; CHECK: DIVERGENT:  i32 %arg6

llvm/test/Analysis/LegacyDivergenceAnalysis/AMDGPU/llvm.amdgcn.buffer.atomic.ll

-;RUN: opt -mtriple=amdgcn-mesa-mesa3d -amdgpu-use-legacy-divergence-analysis -analyze -divergence %s | FileCheck %s
+;RUN: opt -mtriple=amdgcn-mesa-mesa3d -amdgpu-use-legacy-divergence-analysis -enable-new-pm=0 -analyze -divergence %s | FileCheck %s
 
 ;CHECK: DIVERGENT: %orig = call i32 @llvm.amdgcn.buffer.atomic.swap.i32(
 define float @buffer_atomic_swap(<4 x i32> inreg %rsrc, i32 inreg %data) #0 {
 main_body:
   %orig = call i32 @llvm.amdgcn.buffer.atomic.swap.i32(i32 %data, <4 x i32> %rsrc, i32 0, i32 0, i1 0)
   %r = bitcast i32 %orig to float
   ret float %r
 }

llvm/test/Analysis/LegacyDivergenceAnalysis/AMDGPU/llvm.amdgcn.image.atomic.ll

-;RUN: opt -mtriple=amdgcn-mesa-mesa3d -amdgpu-use-legacy-divergence-analysis -analyze -divergence %s | FileCheck %s
+;RUN: opt -mtriple=amdgcn-mesa-mesa3d -amdgpu-use-legacy-divergence-analysis -enable-new-pm=0 -analyze -divergence %s | FileCheck %s
 
 ;CHECK: DIVERGENT: %orig = call i32 @llvm.amdgcn.image.atomic.swap.1d.i32.i32(
 define float @image_atomic_swap(<8 x i32> inreg %rsrc, i32 inreg %addr, i32 inreg %data) #0 {
 main_body:
   %orig = call i32 @llvm.amdgcn.image.atomic.swap.1d.i32.i32(i32 %data, i32 %addr, <8 x i32> %rsrc, i32 0, i32 0)
   %r = bitcast i32 %orig to float
   ret float %r
 }

llvm/test/Analysis/LegacyDivergenceAnalysis/AMDGPU/loads.ll

-; RUN: opt -mtriple=amdgcn-- -amdgpu-use-legacy-divergence-analysis -analyze -divergence %s | FileCheck %s
+; RUN: opt -mtriple=amdgcn-- -amdgpu-use-legacy-divergence-analysis -enable-new-pm=0 -analyze -divergence %s | FileCheck %s
 
 ; Test that we consider loads from flat and private addrspaces to be divergent.
 
 ; CHECK: DIVERGENT: %val = load i32, i32* %flat, align 4
 define amdgpu_kernel void @flat_load(i32* %flat) {
   %val = load i32, i32* %flat, align 4
   ret void
 }
 
 ; CHECK: DIVERGENT: %val = load i32, i32 addrspace(5)* %priv, align 4
 define amdgpu_kernel void @private_load(i32 addrspace(5)* %priv) {
   %val = load i32, i32 addrspace(5)* %priv, align 4
   ret void
 }

llvm/test/Analysis/LegacyDivergenceAnalysis/AMDGPU/no-return-blocks.ll

-; RUN: opt %s -mtriple amdgcn-- -amdgpu-use-legacy-divergence-analysis -analyze -divergence | FileCheck %s
+; RUN: opt %s -mtriple amdgcn-- -amdgpu-use-legacy-divergence-analysis -enable-new-pm=0 -analyze -divergence | FileCheck %s
 
 ; CHECK: DIVERGENT:  %tmp5 = getelementptr inbounds float, float addrspace(1)* %arg, i64 %tmp2
 ; CHECK: DIVERGENT:  %tmp10 = load volatile float, float addrspace(1)* %tmp5, align 4
 ; CHECK: DIVERGENT:  %tmp11 = load volatile float, float addrspace(1)* %tmp5, align 4
 
 ; The post dominator tree does not have a root node in this case
 define amdgpu_kernel void @no_return_blocks(float addrspace(1)* noalias nocapture readonly %arg, float addrspace(1)* noalias nocapture readonly %arg1) #0 {
 bb0:

llvm/test/Analysis/LegacyDivergenceAnalysis/AMDGPU/phi-undef.ll

-; RUN: opt -mtriple=amdgcn-- -amdgpu-use-legacy-divergence-analysis -analyze -divergence %s | FileCheck %s
+; RUN: opt -mtriple=amdgcn-- -amdgpu-use-legacy-divergence-analysis -enable-new-pm=0 -analyze -divergence %s | FileCheck %s
 
 ; CHECK-LABEL: 'test1':
 ; CHECK-NEXT: DIVERGENT: i32 %bound
 ; CHECK: {{^  *}}%counter =
 ; CHECK-NEXT: DIVERGENT: %break = icmp sge i32 %counter, %bound
 ; CHECK-NEXT: DIVERGENT: br i1 %break, label %footer, label %body
 ; CHECK: {{^  *}}%counter.next =
 ; CHECK: {{^  *}}%counter.footer =

llvm/test/Analysis/LegacyDivergenceAnalysis/AMDGPU/unreachable-loop-block.ll

-; RUN: opt %s -mtriple amdgcn-- -amdgpu-use-legacy-divergence-analysis -analyze -divergence | FileCheck %s
+; RUN: opt %s -mtriple amdgcn-- -amdgpu-use-legacy-divergence-analysis -enable-new-pm=0 -analyze -divergence | FileCheck %s
 
 ; CHECK: DIVERGENT:  %tmp = cmpxchg volatile
 define amdgpu_kernel void @unreachable_loop(i32 %tidx) #0 {
 entry:
   unreachable
 
 unreachable_loop:                                        ; preds = %do.body.i, %if.then11
   %tmp = cmpxchg volatile i32 addrspace(1)* null, i32 0, i32 0 seq_cst seq_cst
   %cmp.i = extractvalue { i32, i1 } %tmp, 1
   br i1 %cmp.i, label %unreachable_loop, label %end
 
 end:                                      ; preds = %do.body.i51, %atomicAdd_g_f.exit
   unreachable
 }
 
 attributes #0 = { norecurse nounwind }

llvm/test/Analysis/LegacyDivergenceAnalysis/AMDGPU/workitem-intrinsics.ll

-; RUN: opt  -mtriple amdgcn-unknown-amdhsa -amdgpu-use-legacy-divergence-analysis -analyze -divergence %s | FileCheck %s
+; RUN: opt  -mtriple amdgcn-unknown-amdhsa -amdgpu-use-legacy-divergence-analysis -enable-new-pm=0 -analyze -divergence %s | FileCheck %s
 
 declare i32 @llvm.amdgcn.workitem.id.x() #0
 declare i32 @llvm.amdgcn.workitem.id.y() #0
 declare i32 @llvm.amdgcn.workitem.id.z() #0
 declare i32 @llvm.amdgcn.mbcnt.lo(i32, i32) #0
 declare i32 @llvm.amdgcn.mbcnt.hi(i32, i32) #0
 
 ; CHECK: DIVERGENT:  %id.x = call i32 @llvm.amdgcn.workitem.id.x()

llvm/test/Analysis/LegacyDivergenceAnalysis/NVPTX/diverge.ll

-; RUN: opt %s -analyze -divergence | FileCheck %s
+; RUN: opt %s -enable-new-pm=0 -analyze -divergence | FileCheck %s
 
 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 'Legacy Divergence Analysis' for function 'no_diverge'
 entry:

llvm/tools/opt/opt.cpp

std::vector<StringRef> PassNamePrefix = {
       "x86-",  "xcore-", "wasm-",    "systemz-", "ppc-",    "nvvm-",   "nvptx-",
       "mips-", "lanai-", "hexagon-", "bpf-",     "avr-",    "thumb2-", "arm-",
       "si-",   "gcn-",   "amdgpu-",  "aarch64-", "amdgcn-", "polly-"};
   std::vector<StringRef> PassNameContain = {"ehprepare"};
   std::vector<StringRef> PassNameExact = {
       "safe-stack",           "cost-model",
       "codegenprepare",       "interleaved-load-combine",
       "unreachableblockelim", "verify-safepoint-ir",
-      "divergence",           "atomic-expand",
+      "atomic-expand",
       "hardware-loops",       "type-promotion",
       "mve-tail-predication", "interleaved-access",
       "global-merge",         "pre-isel-intrinsic-lowering",
       "expand-reductions",    "indirectbr-expand",
       "generic-to-nvvm",      "expandmemcmp",
       "loop-reduce",          "lower-amx-type",
       "polyhedral-info",      "replace-with-veclib"};
   for (const auto &P : PassNamePrefix)

llvm/unittests/Analysis/DivergenceAnalysisTest.cpp

BasicBlock *GetBlockByName(StringRef BlockName, Function &F) {
   for (auto &BB : F) {
     if (BB.getName() != BlockName)
       continue;
     return &BB;
   }
   return nullptr;
 }
 
-// We use this fixture to ensure that we clean up DivergenceAnalysis before
+// We use this fixture to ensure that we clean up DivergenceAnalysisImpl before
 // deleting the PassManager.
 class DivergenceAnalysisTest : public testing::Test {
 protected:
   LLVMContext Context;
   Module M;
   TargetLibraryInfoImpl TLII;
   TargetLibraryInfo TLI;
 
   std::unique_ptr<DominatorTree> DT;
   std::unique_ptr<PostDominatorTree> PDT;
   std::unique_ptr<LoopInfo> LI;
   std::unique_ptr<SyncDependenceAnalysis> SDA;
 
   DivergenceAnalysisTest() : M("", Context), TLII(), TLI(TLII) {}
 
-  DivergenceAnalysis buildDA(Function &F, bool IsLCSSA) {
+  DivergenceAnalysisImpl buildDA(Function &F, bool IsLCSSA) {
     DT.reset(new DominatorTree(F));
     PDT.reset(new PostDominatorTree(F));
     LI.reset(new LoopInfo(*DT));
     SDA.reset(new SyncDependenceAnalysis(*DT, *PDT, *LI));
-    return DivergenceAnalysis(F, nullptr, *DT, *LI, *SDA, IsLCSSA);
+    return DivergenceAnalysisImpl(F, nullptr, *DT, *LI, *SDA, IsLCSSA);
   }
 
   void runWithDA(
       Module &M, StringRef FuncName, bool IsLCSSA,
-      function_ref<void(Function &F, LoopInfo &LI, DivergenceAnalysis &DA)>
+      function_ref<void(Function &F, LoopInfo &LI, DivergenceAnalysisImpl &DA)>
           Test) {
     auto *F = M.getFunction(FuncName);
     ASSERT_NE(F, nullptr) << "Could not find " << FuncName;
-    DivergenceAnalysis DA = buildDA(*F, IsLCSSA);
+    DivergenceAnalysisImpl DA = buildDA(*F, IsLCSSA);
     Test(*F, *LI, DA);
   }
 };
 
 // Simple initial state test
 TEST_F(DivergenceAnalysisTest, DAInitialState) {
   IntegerType *IntTy = IntegerType::getInt32Ty(Context);
   FunctionType *FTy =
       FunctionType::get(Type::getVoidTy(Context), {IntTy}, false);
   Function *F = Function::Create(FTy, Function::ExternalLinkage, "f", M);
   BasicBlock *BB = BasicBlock::Create(Context, "entry", F);
   ReturnInst::Create(Context, nullptr, BB);
 
-  DivergenceAnalysis DA = buildDA(*F, false);
+  DivergenceAnalysisImpl DA = buildDA(*F, false);
 
   // Whole function region
   EXPECT_EQ(DA.getRegionLoop(), nullptr);
 
   // No divergence in initial state
   EXPECT_FALSE(DA.hasDetectedDivergence());
 
   // No spurious divergence
std::unique_ptr<Module> M = parseAssemblyString(
       "  br i1 %cond.cont, label %loop, label %for.end.loopexit "
       " "
       "for.end.loopexit: "
       "  ret i32 %iv0 "
       "} ",
       Err, C);
 
   Function *F = M->getFunction("f_1");
-  DivergenceAnalysis DA = buildDA(*F, false);
+  DivergenceAnalysisImpl DA = buildDA(*F, false);
   EXPECT_FALSE(DA.hasDetectedDivergence());
 
   auto ItArg = F->arg_begin();
   ItArg++;
   auto &NArg = *ItArg;
 
   // Seed divergence in argument %n
   DA.markDivergent(NArg);
std::unique_ptr<Module> M = parseAssemblyString(
       "  br label %detached.return "
       " "
       "detached.return: "
       "  ret i32 %val.ret "
       "} ",
       Err, C);
 
   Function *F = M->getFunction("f_lcssa");
-  DivergenceAnalysis DA = buildDA(*F, true);
+  DivergenceAnalysisImpl DA = buildDA(*F, true);
   EXPECT_FALSE(DA.hasDetectedDivergence());
 
   auto ItArg = F->arg_begin();
   ItArg++;
   auto &NArg = *ItArg;
 
   // Seed divergence in argument %n
   DA.markDivergent(NArg);