[Analysis] Remove DA & LegacyDA

[Pierre-vh]

Summary

UniformityAnalysis offers all of the same features and much more, there is no reason left to use the legacy DAs.
See RFC: https://discourse.llvm.org/t/rfc-deprecate-divergenceanalysis-legacydivergenceanalysis/69538

• Remove LegacyDivergenceAnalysis.h/.cpp
• Remove DivergenceAnalysis.h/.cpp + Unit tests
• Remove SyncDependenceAnalysis - it was not a real registered analysis and was only used by DAs
• Remove/adjust references to the passes in the docs where applicable
• Remove TTI hook associated with those passes.
• Move tests to UniformityAnalysis folder.
  • Remove RUN lines for the DA, leave only the UA ones.
• Some tests had to be adjusted/removed depending on how they used the legacy DAs.

[foad]

LGTM with nits.

llvm/docs/ConvergenceAndUniformity.rst

54

Maybe remove the whole sentence?

llvm/include/llvm/ADT/GenericUniformityInfo.h

16

Might as well remove all the commented-out #includes.

llvm/test/CodeGen/AMDGPU/always-uniform.ll

7

What was the point of this part of the test? Should it be updated to test UA instead of LDA?

[Pierre-vh]

llvm/docs/ConvergenceAndUniformity.rst

54

There is a citation in there so I would like to find a way to keep it in or rephrase it
@sameerds you wrote this document, how would you like me to update it?

llvm/test/CodeGen/AMDGPU/always-uniform.ll

7

I'm really not sure what was the point, UA doesn't print anything. Maybe it was just checking DA could run fine on the test without crashing? (though that's already verified by running llc)
Note that I tried using -passes=uniform and it doesn't recognize it, I'm not sure how to run uniformity analysis alone using opt

[Pierre-vh]

Remove includes

[foad]

llvm/lib/Analysis/SyncDependenceAnalysis.cpp

15

I don't think this is true. I think UA uses its own reimplementation of SDA, so this implementation could be removed.

[sameerds]

llvm/docs/ConvergenceAndUniformity.rst

54

You could say "extends previous work on divergence analysis".

llvm/lib/Analysis/SyncDependenceAnalysis.cpp

15

Yeah, nice catch!

[Pierre-vh]

llvm/lib/Analysis/SyncDependenceAnalysis.cpp

15

It's a different analysis so if we remove it, I'd rather make a different patch if that's okay.
Do I just remove this sentence instead?

[foad]

llvm/lib/Analysis/SyncDependenceAnalysis.cpp

15

Sure, remove it or change "is" to "was"...

[sameerds]

llvm/lib/Analysis/SyncDependenceAnalysis.cpp

15

It's not useful on its own. It's not even an LLVM "Analysis" per se, and it was committed along with DA. It's just a CPP file that contains code that does not need to be in the main DA. I think it should entirely removed with DA.

[Pierre-vh]

Remove SDA

[Pierre-vh]

Can someone else also approve the change? For this type of change I'm more comfortable if 2+ people approve before landing :)

llvm/docs/ConvergenceAndUniformity.rst

- These targets must "linearize" a divergent branch to ensure that
   each side of the branch is followed by the corresponding threads in
   the same group. But linearization is unnecessary at uniform
   branches, since the whole group of threads follows either one side
   of the branch or the other.
 
 This document presents a definition of convergence that is reasonable
 for real targets and is compatible with the currently implicit
 semantics of convergent operations in LLVM IR. This is accompanied by
-a *uniformity analysis* that extends the existing divergence analysis
+a *uniformity analysis* that extends the (now removed) divergence analysis
 [DivergenceSPMD]_ to cover irreducible control-flow.
 
 .. [DivergenceSPMD] Julian Rosemann, Simon Moll, and Sebastian
    Hack. 2021. An Abstract Interpretation for SPMD Divergence on
    Reducible Control Flow Graphs. Proc. ACM Program. Lang. 5, POPL,
    Article 31 (January 2021), 35 pages.
    https://doi.org/10.1145/3434312
 

llvm/include/llvm/ADT/GenericUniformityInfo.h

 //===- GenericUniformityInfo.h ---------------------------*- C++ -*--------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef LLVM_ADT_GENERICUNIFORMITYINFO_H
 #define LLVM_ADT_GENERICUNIFORMITYINFO_H
 
 // #include "llvm/ADT/DenseSet.h"
 #include "llvm/ADT/GenericCycleInfo.h"
 // #include "llvm/ADT/SmallPtrSet.h"
 // #include "llvm/ADT/Uniformity.h"
-// #include "llvm/Analysis/LegacyDivergenceAnalysis.h"
 #include "llvm/Support/raw_ostream.h"
 
 namespace llvm {
 
 class TargetTransformInfo;
 
 template <typename ContextT> class GenericUniformityAnalysisImpl;
 template <typename ImplT> struct GenericUniformityAnalysisImplDeleter {

llvm/include/llvm/Analysis/DivergenceAnalysis.h

-//===- llvm/Analysis/DivergenceAnalysis.h - Divergence Analysis -*- C++ -*-===//
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-//===----------------------------------------------------------------------===//
-//
-// \file
-// The divergence analysis determines which instructions and branches are
-// divergent given a set of divergent source instructions.
-//
-//===----------------------------------------------------------------------===//
-
-#ifndef LLVM_ANALYSIS_DIVERGENCEANALYSIS_H
-#define LLVM_ANALYSIS_DIVERGENCEANALYSIS_H
-
-#include "llvm/ADT/DenseSet.h"
-#include "llvm/Analysis/SyncDependenceAnalysis.h"
-#include "llvm/IR/PassManager.h"
-#include <vector>
-
-namespace llvm {
-class Function;
-class Instruction;
-class Loop;
-class raw_ostream;
-class TargetTransformInfo;
-class Value;
-
-/// \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 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 LCSSA form.
-  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.
-  bool inRegion(const Instruction &I) const;
-
-  /// \brief Mark \p UniVal as a value that is always uniform.
-  void addUniformOverride(const Value &UniVal);
-
-  /// \brief Mark \p DivVal as a value that is always divergent. Will not do so
-  /// if `isAlwaysUniform(DivVal)`.
-  /// \returns Whether the tracked divergence state of \p DivVal changed.
-  bool markDivergent(const Value &DivVal);
-
-  /// \brief Propagate divergence to all instructions in the region.
-  /// Divergence is seeded by calls to \p markDivergent.
-  void compute();
-
-  /// \brief Whether any value was marked or analyzed to be divergent.
-  bool hasDetectedDivergence() const { return !DivergentValues.empty(); }
-
-  /// \brief Whether \p Val will always return a uniform value regardless of its
-  /// operands
-  bool isAlwaysUniform(const Value &Val) const;
-
-  /// \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;
-
-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);
-
-  /// \brief Identify all Instructions that become divergent because \p DivExit
-  /// is a divergent loop exit of \p DivLoop. Mark those instructions as
-  /// divergent and push them on the worklist.
-  void propagateLoopExitDivergence(const BasicBlock &DivExit,
-                                   const Loop &DivLoop);
-
-  /// \brief Internal implementation function for propagateLoopExitDivergence.
-  void analyzeLoopExitDivergence(const BasicBlock &DivExit,
-                                 const Loop &OuterDivLoop);
-
-  /// \brief Mark all instruction as divergent that use a value defined in \p
-  /// OuterDivLoop. Push their users on the worklist.
-  void analyzeTemporalDivergence(const Instruction &I,
-                                 const Loop &OuterDivLoop);
-
-  /// \brief Push all users of \p Val (in the region) to the worklist.
-  void pushUsers(const Value &I);
-
-  /// \brief Whether \p Val is divergent when read in \p ObservingBlock.
-  bool isTemporalDivergent(const BasicBlock &ObservingBlock,
-                           const Value &Val) const;
-
-private:
-  const Function &F;
-  // If regionLoop != nullptr, analysis is only performed within \p RegionLoop.
-  // Otherwise, analyze the whole function
-  const Loop *RegionLoop;
-
-  const DominatorTree &DT;
-  const LoopInfo &LI;
-
-  // Recognized divergent loops
-  DenseSet<const Loop *> DivergentLoops;
-
-  // The SDA links divergent branches to divergent control-flow joins.
-  SyncDependenceAnalysis &SDA;
-
-  // Use simplified code path for LCSSA form.
-  bool IsLCSSAForm;
-
-  // Set of known-uniform values.
-  DenseSet<const Value *> UniformOverrides;
-
-  // Detected/marked divergent values.
-  DenseSet<const Value *> DivergentValues;
-
-  // Internal worklist for divergence propagation.
-  std::vector<const Instruction *> Worklist;
-};
-
-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 = false;
-  std::unique_ptr<SyncDependenceAnalysis> SDA;
-  std::unique_ptr<DivergenceAnalysisImpl> DA;
-
-public:
-  DivergenceInfo(Function &F, const DominatorTree &DT,
-                 const PostDominatorTree &PDT, const LoopInfo &LI,
-                 const TargetTransformInfo &TTI, bool KnownReducible);
-
-  /// Whether any divergence was detected.
-  bool hasDivergence() const {
-    return ContainsIrreducible || DA->hasDetectedDivergence();
-  }
-
-  /// The GPU kernel this analysis result is for
-  const Function &getFunction() const { return F; }
-
-  /// Whether \p V is divergent at its definition.
-  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 {
-    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); }
-};
-
-/// \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

-//===- llvm/Analysis/LegacyDivergenceAnalysis.h - KernelDivergence Analysis -*- C++ -*-===//
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-//===----------------------------------------------------------------------===//
-//
-// The kernel divergence analysis is an LLVM pass which can be used to find out
-// if a branch instruction in a GPU program (kernel) is divergent or not. It can help
-// branch optimizations such as jump threading and loop unswitching to make
-// better decisions.
-//
-//===----------------------------------------------------------------------===//
-#ifndef LLVM_ANALYSIS_LEGACYDIVERGENCEANALYSIS_H
-#define LLVM_ANALYSIS_LEGACYDIVERGENCEANALYSIS_H
-
-#include "llvm/ADT/DenseSet.h"
-#include "llvm/Analysis/LoopInfo.h"
-#include "llvm/Analysis/PostDominators.h"
-#include "llvm/IR/PassManager.h"
-#include "llvm/Pass.h"
-#include <memory>
-
-namespace llvm {
-class DivergenceInfo;
-class Function;
-class Module;
-class raw_ostream;
-class TargetTransformInfo;
-class Use;
-class Value;
-
-class LegacyDivergenceAnalysisImpl {
-public:
-  // Returns true if V is divergent at its definition.
-  bool isDivergent(const Value *V) const;
-
-  // Returns true if U is divergent. Uses of a uniform value can be divergent.
-  bool isDivergentUse(const Use *U) const;
-
-  // Returns true if V is uniform/non-divergent.
-  bool isUniform(const Value *V) const { return !isDivergent(V); }
-
-  // Returns true if U is uniform/non-divergent. Uses of a uniform value can be
-  // divergent.
-  bool isUniformUse(const Use *U) const { return !isDivergentUse(U); }
-
-  // Keep the analysis results uptodate by removing an erased value.
-  void removeValue(const Value *V) { DivergentValues.erase(V); }
-
-  // Print all divergent branches in the function.
-  void print(raw_ostream &OS, const Module *) const;
-
-  // Whether analysis should be performed by GPUDivergenceAnalysis.
-  bool shouldUseGPUDivergenceAnalysis(const Function &F,
-                                      const TargetTransformInfo &TTI,
-                                      const LoopInfo &LI);
-
-  void run(Function &F, TargetTransformInfo &TTI, DominatorTree &DT,
-           PostDominatorTree &PDT, const LoopInfo &LI);
-
-protected:
-  // (optional) handle to new DivergenceAnalysis
-  std::unique_ptr<DivergenceInfo> gpuDA;
-
-  // Stores all divergent values.
-  DenseSet<const Value *> DivergentValues;
-
-  // Stores divergent uses of possibly uniform values.
-  DenseSet<const Use *> DivergentUses;
-};
-
-class LegacyDivergenceAnalysis : public FunctionPass,
-                                 public LegacyDivergenceAnalysisImpl {
-public:
-  static char ID;
-
-  LegacyDivergenceAnalysis();
-  void getAnalysisUsage(AnalysisUsage &AU) const override;
-  bool runOnFunction(Function &F) override;
-};
-
-class LegacyDivergenceAnalysisPass
-    : public PassInfoMixin<LegacyDivergenceAnalysisPass>,
-      public LegacyDivergenceAnalysisImpl {
-public:
-  PreservedAnalyses run(Function &F, FunctionAnalysisManager &AM);
-
-private:
-  // (optional) handle to new DivergenceAnalysis
-  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 namespace llvm
-
-#endif // LLVM_ANALYSIS_LEGACYDIVERGENCEANALYSIS_H

llvm/include/llvm/Analysis/Passes.h

namespace llvm {
   //
   // createDelinearizationPass - This pass implements attempts to restore
   // multidimensional array indices from linearized expressions.
   //
   FunctionPass *createDelinearizationPass();
 
   //===--------------------------------------------------------------------===//
   //
-  // createLegacyDivergenceAnalysisPass - This pass determines which branches in a GPU
-  // program are divergent.
-  //
-  FunctionPass *createLegacyDivergenceAnalysisPass();
-
-  //===--------------------------------------------------------------------===//
-  //
   // 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.

llvm/include/llvm/Analysis/TargetTransformInfo.h

public:
 
   /// 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;
 
-  /// Return true if the target prefers to use GPU divergence analysis to
-  /// replace the legacy version.
-  bool useGPUDivergenceAnalysis() const;
-
   /// Returns whether V is a source of divergence.
   ///
   /// This function provides the target-dependent information for
-  /// the target-independent LegacyDivergenceAnalysis. LegacyDivergenceAnalysis
-  /// first builds the dependency graph, and then runs the reachability
-  /// algorithm starting with the sources of divergence.
+  /// the target-independent UniformityAnalysis.
   bool isSourceOfDivergence(const Value *V) const;
 
   // Returns true for the target specific
   // set of operations which produce uniform result
   // even taking non-uniform arguments
   bool isAlwaysUniform(const Value *V) const;
 
   /// Query the target whether the specified address space cast from FromAS to
public:
   getEstimatedNumberOfCaseClusters(const SwitchInst &SI, unsigned &JTSize,
                                    ProfileSummaryInfo *PSI,
                                    BlockFrequencyInfo *BFI) = 0;
   virtual InstructionCost getInstructionCost(const User *U,
                                              ArrayRef<const Value *> Operands,
                                              TargetCostKind CostKind) = 0;
   virtual BranchProbability getPredictableBranchThreshold() = 0;
   virtual bool hasBranchDivergence() = 0;
-  virtual bool useGPUDivergenceAnalysis() = 0;
   virtual bool isSourceOfDivergence(const Value *V) = 0;
   virtual bool isAlwaysUniform(const Value *V) = 0;
   virtual bool isValidAddrSpaceCast(unsigned FromAS, unsigned ToAS) const = 0;
   virtual unsigned getFlatAddressSpace() = 0;
   virtual bool collectFlatAddressOperands(SmallVectorImpl<int> &OpIndexes,
                                           Intrinsic::ID IID) const = 0;
   virtual bool isNoopAddrSpaceCast(unsigned FromAS, unsigned ToAS) const = 0;
   virtual bool
InstructionCost getInstructionCost(const User *U,
                                      ArrayRef<const Value *> Operands,
                                      TargetCostKind CostKind) override {
     return Impl.getInstructionCost(U, Operands, CostKind);
   }
   BranchProbability getPredictableBranchThreshold() override {
     return Impl.getPredictableBranchThreshold();
   }
   bool hasBranchDivergence() override { return Impl.hasBranchDivergence(); }
-  bool useGPUDivergenceAnalysis() override {
-    return Impl.useGPUDivergenceAnalysis();
-  }
   bool isSourceOfDivergence(const Value *V) override {
     return Impl.isSourceOfDivergence(V);
   }
 
   bool isAlwaysUniform(const Value *V) override {
     return Impl.isAlwaysUniform(V);
   }
 

llvm/include/llvm/Analysis/TargetTransformInfoImpl.h

public:
   // the loss from (100 - N)% mispredicted executions for the majority of
   //  intended targets.
   BranchProbability getPredictableBranchThreshold() const {
     return BranchProbability(99, 100);
   }
 
   bool hasBranchDivergence() const { return false; }
 
-  bool useGPUDivergenceAnalysis() const { return false; }
-
   bool isSourceOfDivergence(const Value *V) const { return false; }
 
   bool isAlwaysUniform(const Value *V) const { return false; }
 
   bool isValidAddrSpaceCast(unsigned FromAS, unsigned ToAS) const {
     return false;
   }
 

llvm/include/llvm/CodeGen/BasicTTIImpl.h

bool allowsMisalignedMemoryAccesses(LLVMContext &Context, unsigned BitWidth,
                                       unsigned *Fast) const {
     EVT E = EVT::getIntegerVT(Context, BitWidth);
     return getTLI()->allowsMisalignedMemoryAccesses(
         E, AddressSpace, Alignment, MachineMemOperand::MONone, Fast);
   }
 
   bool hasBranchDivergence() { return false; }
 
-  bool useGPUDivergenceAnalysis() { return false; }
-
   bool isSourceOfDivergence(const Value *V) { return false; }
 
   bool isAlwaysUniform(const Value *V) { return false; }
 
   bool isValidAddrSpaceCast(unsigned FromAS, unsigned ToAS) const {
     return false;
   }
 

llvm/include/llvm/InitializePasses.h

 void initializeJumpThreadingPass(PassRegistry&);
 void initializeLCSSAVerificationPassPass(PassRegistry&);
 void initializeLCSSAWrapperPassPass(PassRegistry&);
 void initializeLazyBlockFrequencyInfoPassPass(PassRegistry&);
 void initializeLazyBranchProbabilityInfoPassPass(PassRegistry&);
 void initializeLazyMachineBlockFrequencyInfoPassPass(PassRegistry&);
 void initializeLazyValueInfoPrinterPass(PassRegistry&);
 void initializeLazyValueInfoWrapperPassPass(PassRegistry&);
-void initializeLegacyDivergenceAnalysisPass(PassRegistry&);
 void initializeLegacyLICMPassPass(PassRegistry&);
 void initializeLegacyLoopSinkPassPass(PassRegistry&);
 void initializeLegalizerPass(PassRegistry&);
 void initializeGISelCSEAnalysisWrapperPassPass(PassRegistry &);
 void initializeGISelKnownBitsAnalysisPass(PassRegistry &);
 void initializeLintLegacyPassPass(PassRegistry &);
 void initializeLiveDebugValuesPass(PassRegistry&);
 void initializeLiveDebugVariablesPass(PassRegistry&);

llvm/include/llvm/LinkAllPasses.h

ForcePassLinking() {
       (void) llvm::createGlobalsAAWrapperPass();
       (void) llvm::createGuardWideningPass();
       (void) llvm::createLoopGuardWideningPass();
       (void) llvm::createInductiveRangeCheckEliminationPass();
       (void) llvm::createInstSimplifyLegacyPass();
       (void) llvm::createInstructionCombiningPass();
       (void) llvm::createJMCInstrumenterPass();
       (void) llvm::createLCSSAPass();
-      (void) llvm::createLegacyDivergenceAnalysisPass();
       (void) llvm::createLICMPass();
       (void) llvm::createLoopSinkPass();
       (void) llvm::createLazyValueInfoPass();
       (void) llvm::createLoopExtractorPass();
       (void) llvm::createLoopPredicationPass();
       (void) llvm::createLoopSimplifyPass();
       (void) llvm::createLoopSimplifyCFGPass();
       (void) llvm::createLoopStrengthReducePass();

llvm/lib/Analysis/Analysis.cpp

void llvm::initializeAnalysis(PassRegistry &Registry) {
   initializeIVUsersWrapperPassPass(Registry);
   initializeInstCountLegacyPassPass(Registry);
   initializeIntervalPartitionPass(Registry);
   initializeIRSimilarityIdentifierWrapperPassPass(Registry);
   initializeLazyBranchProbabilityInfoPassPass(Registry);
   initializeLazyBlockFrequencyInfoPassPass(Registry);
   initializeLazyValueInfoWrapperPassPass(Registry);
   initializeLazyValueInfoPrinterPass(Registry);
-  initializeLegacyDivergenceAnalysisPass(Registry);
   initializeLintLegacyPassPass(Registry);
   initializeLoopInfoWrapperPassPass(Registry);
   initializeMemDepPrinterPass(Registry);
   initializeMemDerefPrinterPass(Registry);
   initializeMemoryDependenceWrapperPassPass(Registry);
   initializeModuleDebugInfoLegacyPrinterPass(Registry);
   initializeModuleSummaryIndexWrapperPassPass(Registry);
   initializeMustExecutePrinterPass(Registry);

llvm/lib/Analysis/CMakeLists.txt

add_llvm_component_library(LLVMAnalysis
   DDG.cpp
   DDGPrinter.cpp
   ConstraintSystem.cpp
   Delinearization.cpp
   DemandedBits.cpp
   DependenceAnalysis.cpp
   DependenceGraphBuilder.cpp
   DevelopmentModeInlineAdvisor.cpp
-  DivergenceAnalysis.cpp
   DomPrinter.cpp
   DomTreeUpdater.cpp
   DominanceFrontier.cpp
   FunctionPropertiesAnalysis.cpp
   GlobalsModRef.cpp
   GuardUtils.cpp
   HeatUtils.cpp
   IRSimilarityIdentifier.cpp
add_llvm_component_library(LLVMAnalysis
   InstructionSimplify.cpp
   InteractiveModelRunner.cpp
   Interval.cpp
   IntervalPartition.cpp
   LazyBranchProbabilityInfo.cpp
   LazyBlockFrequencyInfo.cpp
   LazyCallGraph.cpp
   LazyValueInfo.cpp
-  LegacyDivergenceAnalysis.cpp
   Lint.cpp
   Loads.cpp
   Local.cpp
   LoopAccessAnalysis.cpp
   LoopAnalysisManager.cpp
   LoopCacheAnalysis.cpp
   LoopNestAnalysis.cpp
   LoopUnrollAnalyzer.cpp

llvm/lib/Analysis/DivergenceAnalysis.cpp

-//===---- DivergenceAnalysis.cpp --- Divergence Analysis Implementation ----==//
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-//===----------------------------------------------------------------------===//
-//
-// This file implements a general divergence analysis for loop vectorization
-// and GPU programs. It determines which branches and values in a loop or GPU
-// program are 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 execution 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 re-converge.
-//
-// Due to this execution model, some optimizations such as jump
-// threading and loop unswitching can interfere with thread re-convergence.
-// Therefore, an analysis that computes which branches in a GPU program are
-// divergent can help the compiler to selectively run these optimizations.
-//
-// This implementation is derived from the Vectorization Analysis of the
-// Region Vectorizer (RV). The analysis is based on the approach described in
-//
-//   An abstract interpretation for SPMD divergence
-//       on reducible control flow graphs.
-//   Julian Rosemann, Simon Moll and Sebastian Hack
-//   POPL '21
-//
-// 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
-// (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,
-//
-//   %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.
-//
-// The sync dependence detection (which branch induces divergence in which join
-// points) is implemented in the SyncDependenceAnalysis.
-//
-// 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/ADT/PostOrderIterator.h"
-#include "llvm/Analysis/CFG.h"
-#include "llvm/Analysis/LoopInfo.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/Value.h"
-#include "llvm/Support/Debug.h"
-#include "llvm/Support/raw_ostream.h"
-
-using namespace llvm;
-
-#define DEBUG_TYPE "divergence"
-
-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 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 DivergenceAnalysisImpl::addUniformOverride(const Value &UniVal) {
-  UniformOverrides.insert(&UniVal);
-}
-
-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 DivergenceAnalysisImpl::inRegion(const Instruction &I) const {
-  return I.getParent() && inRegion(*I.getParent());
-}
-
-bool DivergenceAnalysisImpl::inRegion(const BasicBlock &BB) const {
-  return RegionLoop ? RegionLoop->contains(&BB) : (BB.getParent() == &F);
-}
-
-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()) {
-    const auto *UserInst = dyn_cast<const Instruction>(User);
-    if (!UserInst)
-      continue;
-
-    // only compute divergent inside loop
-    if (!inRegion(*UserInst))
-      continue;
-
-    // All users of divergent values are immediate divergent
-    if (markDivergent(*UserInst))
-      Worklist.push_back(UserInst);
-  }
-}
-
-static const Instruction *getIfCarriedInstruction(const Use &U,
-                                                  const Loop &DivLoop) {
-  const auto *I = dyn_cast<const Instruction>(&U);
-  if (!I)
-    return nullptr;
-  if (!DivLoop.contains(I))
-    return nullptr;
-  return I;
-}
-
-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 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 non-LCSSA we have to follow all live out edges wherever they may lead.
-  const BasicBlock &LoopHeader = *OuterDivLoop.getHeader();
-  SmallVector<const BasicBlock *, 8> TaintStack;
-  TaintStack.push_back(&DivExit);
-
-  // Otherwise potential users of loop-carried values could be anywhere in the
-  // dominance region of DivLoop (including its fringes for phi nodes)
-  DenseSet<const BasicBlock *> Visited;
-  Visited.insert(&DivExit);
-
-  do {
-    auto *UserBlock = TaintStack.pop_back_val();
-
-    // don't spread divergence beyond the region
-    if (!inRegion(*UserBlock))
-      continue;
-
-    assert(!OuterDivLoop.contains(UserBlock) &&
-           "irreducible control flow detected");
-
-    // phi nodes at the fringes of the dominance region
-    if (!DT.dominates(&LoopHeader, UserBlock)) {
-      // all PHI nodes of UserBlock become divergent
-      for (const auto &Phi : UserBlock->phis()) {
-        analyzeTemporalDivergence(Phi, OuterDivLoop);
-      }
-      continue;
-    }
-
-    // Taint outside users of values carried by OuterDivLoop.
-    for (const auto &I : *UserBlock) {
-      analyzeTemporalDivergence(I, OuterDivLoop);
-    }
-
-    // visit all blocks in the dominance region
-    for (const auto *SuccBlock : successors(UserBlock)) {
-      if (!Visited.insert(SuccBlock).second) {
-        continue;
-      }
-      TaintStack.push_back(SuccBlock);
-    }
-  } while (!TaintStack.empty());
-}
-
-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 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 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 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 DivergenceAnalysisImpl::isAlwaysUniform(const Value &V) const {
-  return UniformOverrides.contains(&V);
-}
-
-bool DivergenceAnalysisImpl::isDivergent(const Value &V) const {
-  return DivergentValues.contains(&V);
-}
-
-bool DivergenceAnalysisImpl::isDivergentUse(const Use &U) const {
-  Value &V = *U.get();
-  Instruction &I = *cast<Instruction>(U.getUser());
-  return isDivergent(V) || isTemporalDivergent(*I.getParent(), V);
-}
-
-DivergenceInfo::DivergenceInfo(Function &F, const DominatorTree &DT,
-                               const PostDominatorTree &PDT, const LoopInfo &LI,
-                               const TargetTransformInfo &TTI,
-                               bool KnownReducible)
-    : F(F) {
-  if (!KnownReducible) {
-    using RPOTraversal = ReversePostOrderTraversal<const Function *>;
-    RPOTraversal FuncRPOT(&F);
-    if (containsIrreducibleCFG<const BasicBlock *, const RPOTraversal,
-                               const LoopInfo>(FuncRPOT, LI)) {
-      ContainsIrreducible = true;
-      return;
-    }
-  }
-  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);
-    } else if (TTI.isAlwaysUniform(&I)) {
-      DA->addUniformOverride(I);
-    }
-  }
-  for (auto &Arg : F.args()) {
-    if (TTI.isSourceOfDivergence(&Arg)) {
-      DA->markDivergent(Arg);
-    }
-  }
-
-  DA->compute();
-}
-
-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);
-
-  return DivergenceInfo(F, DT, PDT, LI, TTI, /* KnownReducible = */ false);
-}
-
-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 (const BasicBlock &BB : F) {
-      OS << "\n           " << BB.getName() << ":\n";
-      for (const auto &I : BB.instructionsWithoutDebug()) {
-        OS << (DI.isDivergent(I) ? "DIVERGENT:     " : "               ");
-        OS << I << "\n";
-      }
-    }
-  }
-  return PreservedAnalyses::all();
-}

llvm/lib/Analysis/LegacyDivergenceAnalysis.cpp

-//===- LegacyDivergenceAnalysis.cpp --------- Legacy Divergence Analysis
-//Implementation -==//
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-//===----------------------------------------------------------------------===//
-//
-// This file implements 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 execution 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 program. It
-// partially implements the approach described in
-//
-//   Divergence Analysis
-//   Sampaio, Souza, Collange, Pereira
-//   TOPLAS '13
-//
-// 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.
-//
-// 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.
-//
-//===----------------------------------------------------------------------===//
-
-#include "llvm/Analysis/LegacyDivergenceAnalysis.h"
-#include "llvm/ADT/PostOrderIterator.h"
-#include "llvm/Analysis/CFG.h"
-#include "llvm/Analysis/DivergenceAnalysis.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/Value.h"
-#include "llvm/InitializePasses.h"
-#include "llvm/Support/CommandLine.h"
-#include "llvm/Support/Debug.h"
-#include "llvm/Support/raw_ostream.h"
-#include <vector>
-using namespace llvm;
-
-#define DEBUG_TYPE "divergence"
-
-// transparently use the GPUDivergenceAnalysis
-static cl::opt<bool> UseGPUDA("use-gpu-divergence-analysis", cl::init(false),
-                              cl::Hidden,
-                              cl::desc("turn the LegacyDivergenceAnalysis into "
-                                       "a wrapper for GPUDivergenceAnalysis"));
-
-namespace {
-
-class DivergencePropagator {
-public:
-  DivergencePropagator(Function &F, TargetTransformInfo &TTI, DominatorTree &DT,
-                       PostDominatorTree &PDT, DenseSet<const Value *> &DV,
-                       DenseSet<const Use *> &DU)
-      : F(F), TTI(TTI), DT(DT), PDT(PDT), DV(DV), DU(DU) {}
-  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 TI.
-  void exploreSyncDependency(Instruction *TI);
-  // Computes the influence region from Start to End. This region includes all
-  // basic blocks on any simple path from Start to End.
-  void computeInfluenceRegion(BasicBlock *Start, BasicBlock *End,
-                              DenseSet<BasicBlock *> &InfluenceRegion);
-  // Finds all users of I that are outside the influence region, and add these
-  // users to Worklist.
-  void findUsersOutsideInfluenceRegion(
-      Instruction &I, const DenseSet<BasicBlock *> &InfluenceRegion);
-
-  Function &F;
-  TargetTransformInfo &TTI;
-  DominatorTree &DT;
-  PostDominatorTree &PDT;
-  std::vector<Value *> Worklist; // Stack for DFS.
-  DenseSet<const Value *> &DV;   // Stores all divergent values.
-  DenseSet<const Use *> &DU;   // Stores divergent uses of possibly uniform
-                               // values.
-};
-
-void DivergencePropagator::populateWithSourcesOfDivergence() {
-  Worklist.clear();
-  DV.clear();
-  DU.clear();
-  for (auto &I : instructions(F)) {
-    if (TTI.isSourceOfDivergence(&I)) {
-      Worklist.push_back(&I);
-      DV.insert(&I);
-    }
-  }
-  for (auto &Arg : F.args()) {
-    if (TTI.isSourceOfDivergence(&Arg)) {
-      Worklist.push_back(&Arg);
-      DV.insert(&Arg);
-    }
-  }
-}
-
-void DivergencePropagator::exploreSyncDependency(Instruction *TI) {
-  // Propagation rule 1: if branch TI is divergent, all PHINodes in TI's
-  // immediate post dominator are divergent. This rule handles if-then-else
-  // patterns. For example,
-  //
-  // if (tid < 5)
-  //   a1 = 1;
-  // else
-  //   a2 = 2;
-  // a = phi(a1, a2); // sync dependent on (tid < 5)
-  BasicBlock *ThisBB = TI->getParent();
-
-  // Unreachable blocks may not be in the dominator tree.
-  if (!DT.isReachableFromEntry(ThisBB))
-    return;
-
-  // If the function has no exit blocks or doesn't reach any exit blocks, the
-  // post dominator may be null.
-  DomTreeNode *ThisNode = PDT.getNode(ThisBB);
-  if (!ThisNode)
-    return;
-
-  BasicBlock *IPostDom = ThisNode->getIDom()->getBlock();
-  if (IPostDom == nullptr)
-    return;
-
-  for (auto I = IPostDom->begin(); isa<PHINode>(I); ++I) {
-    // A PHINode is uniform if it returns the same value no matter which path is
-    // taken.
-    if (!cast<PHINode>(I)->hasConstantOrUndefValue() && DV.insert(&*I).second)
-      Worklist.push_back(&*I);
-  }
-
-  // Propagation rule 2: if a value defined in a loop is used outside, the user
-  // is sync dependent on the condition of the loop exits that dominate the
-  // user. For example,
-  //
-  // int i = 0;
-  // do {
-  //   i++;
-  //   if (foo(i)) ... // uniform
-  // } while (i < tid);
-  // if (bar(i)) ...   // divergent
-  //
-  // A program may contain unstructured loops. Therefore, we cannot leverage
-  // LoopInfo, which only recognizes natural loops.
-  //
-  // The algorithm used here handles both natural and unstructured loops.  Given
-  // a branch TI, we first compute its influence region, the union of all simple
-  // paths from TI to its immediate post dominator (IPostDom). Then, we search
-  // for all the values defined in the influence region but used outside. All
-  // these users are sync dependent on TI.
-  DenseSet<BasicBlock *> InfluenceRegion;
-  computeInfluenceRegion(ThisBB, IPostDom, InfluenceRegion);
-  // An insight that can speed up the search process is that all the in-region
-  // values that are used outside must dominate TI. Therefore, instead of
-  // searching every basic blocks in the influence region, we search all the
-  // dominators of TI until it is outside the influence region.
-  BasicBlock *InfluencedBB = ThisBB;
-  while (InfluenceRegion.count(InfluencedBB)) {
-    for (auto &I : *InfluencedBB) {
-      if (!DV.count(&I))
-        findUsersOutsideInfluenceRegion(I, InfluenceRegion);
-    }
-    DomTreeNode *IDomNode = DT.getNode(InfluencedBB)->getIDom();
-    if (IDomNode == nullptr)
-      break;
-    InfluencedBB = IDomNode->getBlock();
-  }
-}
-
-void DivergencePropagator::findUsersOutsideInfluenceRegion(
-    Instruction &I, const DenseSet<BasicBlock *> &InfluenceRegion) {
-  for (Use &Use : I.uses()) {
-    Instruction *UserInst = cast<Instruction>(Use.getUser());
-    if (!InfluenceRegion.count(UserInst->getParent())) {
-      DU.insert(&Use);
-      if (DV.insert(UserInst).second)
-        Worklist.push_back(UserInst);
-    }
-  }
-}
-
-// A helper function for computeInfluenceRegion that adds successors of "ThisBB"
-// to the influence region.
-static void
-addSuccessorsToInfluenceRegion(BasicBlock *ThisBB, BasicBlock *End,
-                               DenseSet<BasicBlock *> &InfluenceRegion,
-                               std::vector<BasicBlock *> &InfluenceStack) {
-  for (BasicBlock *Succ : successors(ThisBB)) {
-    if (Succ != End && InfluenceRegion.insert(Succ).second)
-      InfluenceStack.push_back(Succ);
-  }
-}
-
-void DivergencePropagator::computeInfluenceRegion(
-    BasicBlock *Start, BasicBlock *End,
-    DenseSet<BasicBlock *> &InfluenceRegion) {
-  assert(PDT.properlyDominates(End, Start) &&
-         "End does not properly dominate Start");
-
-  // The influence region starts from the end of "Start" to the beginning of
-  // "End". Therefore, "Start" should not be in the region unless "Start" is in
-  // a loop that doesn't contain "End".
-  std::vector<BasicBlock *> InfluenceStack;
-  addSuccessorsToInfluenceRegion(Start, End, InfluenceRegion, InfluenceStack);
-  while (!InfluenceStack.empty()) {
-    BasicBlock *BB = InfluenceStack.back();
-    InfluenceStack.pop_back();
-    addSuccessorsToInfluenceRegion(BB, End, InfluenceRegion, InfluenceStack);
-  }
-}
-
-void DivergencePropagator::exploreDataDependency(Value *V) {
-  // Follow def-use chains of V.
-  for (User *U : V->users()) {
-    if (!TTI.isAlwaysUniform(U) && DV.insert(U).second)
-      Worklist.push_back(U);
-  }
-}
-
-void DivergencePropagator::propagate() {
-  // Traverse the dependency graph using DFS.
-  while (!Worklist.empty()) {
-    Value *V = Worklist.back();
-    Worklist.pop_back();
-    if (Instruction *I = dyn_cast<Instruction>(V)) {
-      // Terminators with less than two successors won't introduce sync
-      // dependency. Ignore them.
-      if (I->isTerminator() && I->getNumSuccessors() > 1)
-        exploreSyncDependency(I);
-    }
-    exploreDataDependency(V);
-  }
-}
-
-} // namespace
-
-// Register this pass.
-char LegacyDivergenceAnalysis::ID = 0;
-LegacyDivergenceAnalysis::LegacyDivergenceAnalysis() : FunctionPass(ID) {
-  initializeLegacyDivergenceAnalysisPass(*PassRegistry::getPassRegistry());
-}
-INITIALIZE_PASS_BEGIN(LegacyDivergenceAnalysis, "divergence",
-                      "Legacy Divergence Analysis", false, true)
-INITIALIZE_PASS_DEPENDENCY(DominatorTreeWrapperPass)
-INITIALIZE_PASS_DEPENDENCY(PostDominatorTreeWrapperPass)
-INITIALIZE_PASS_DEPENDENCY(LoopInfoWrapperPass)
-INITIALIZE_PASS_END(LegacyDivergenceAnalysis, "divergence",
-                    "Legacy Divergence Analysis", false, true)
-
-FunctionPass *llvm::createLegacyDivergenceAnalysisPass() {
-  return new LegacyDivergenceAnalysis();
-}
-
-bool LegacyDivergenceAnalysisImpl::shouldUseGPUDivergenceAnalysis(
-    const Function &F, const TargetTransformInfo &TTI, const LoopInfo &LI) {
-  if (!(UseGPUDA || TTI.useGPUDivergenceAnalysis()))
-    return false;
-
-  // GPUDivergenceAnalysis requires a reducible CFG.
-  using RPOTraversal = ReversePostOrderTraversal<const Function *>;
-  RPOTraversal FuncRPOT(&F);
-  return !containsIrreducibleCFG<const BasicBlock *, const RPOTraversal,
-                                 const LoopInfo>(FuncRPOT, LI);
-}
-
-void LegacyDivergenceAnalysisImpl::run(Function &F,
-                                       llvm::TargetTransformInfo &TTI,
-                                       llvm::DominatorTree &DT,
-                                       llvm::PostDominatorTree &PDT,
-                                       const llvm::LoopInfo &LI) {
-  if (shouldUseGPUDivergenceAnalysis(F, TTI, LI)) {
-    // run the new GPU divergence analysis
-    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();
-  }
-}
-
-bool LegacyDivergenceAnalysisImpl::isDivergent(const Value *V) const {
-  if (gpuDA) {
-    return gpuDA->isDivergent(*V);
-  }
-  return DivergentValues.count(V);
-}
-
-bool LegacyDivergenceAnalysisImpl::isDivergentUse(const Use *U) const {
-  if (gpuDA) {
-    return gpuDA->isDivergentUse(*U);
-  }
-  return DivergentValues.count(U->get()) || DivergentUses.count(U);
-}
-
-void LegacyDivergenceAnalysisImpl::print(raw_ostream &OS,
-                                         const Module *) const {
-  if ((!gpuDA || !gpuDA->hasDivergence()) && DivergentValues.empty())
-    return;
-
-  const Function *F = nullptr;
-  if (!DivergentValues.empty()) {
-    const Value *FirstDivergentValue = *DivergentValues.begin();
-    if (const Argument *Arg = dyn_cast<Argument>(FirstDivergentValue)) {
-      F = Arg->getParent();
-    } else if (const Instruction *I =
-                   dyn_cast<Instruction>(FirstDivergentValue)) {
-      F = I->getParent()->getParent();
-    } else {
-      llvm_unreachable("Only arguments and instructions can be divergent");
-    }
-  } else if (gpuDA) {
-    F = &gpuDA->getFunction();
-  }
-  if (!F)
-    return;
-
-  // Dumps all divergent values in F, arguments and then instructions.
-  for (const auto &Arg : F->args()) {
-    OS << (isDivergent(&Arg) ? "DIVERGENT: " : "           ");
-    OS << Arg << "\n";
-  }
-  // Iterate instructions using instructions() to ensure a deterministic order.
-  for (const BasicBlock &BB : *F) {
-    OS << "\n           " << BB.getName() << ":\n";
-    for (const auto &I : BB.instructionsWithoutDebug()) {
-      OS << (isDivergent(&I) ? "DIVERGENT:     " : "               ");
-      OS << I << "\n";
-    }
-  }
-  OS << "\n";
-}
-
-void LegacyDivergenceAnalysis::getAnalysisUsage(AnalysisUsage &AU) const {
-  AU.addRequiredTransitive<DominatorTreeWrapperPass>();
-  AU.addRequiredTransitive<PostDominatorTreeWrapperPass>();
-  AU.addRequiredTransitive<LoopInfoWrapperPass>();
-  AU.setPreservesAll();
-}
-
-bool LegacyDivergenceAnalysis::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;
-
-  DivergentValues.clear();
-  DivergentUses.clear();
-  gpuDA = nullptr;
-
-  auto &DT = getAnalysis<DominatorTreeWrapperPass>().getDomTree();
-  auto &PDT = getAnalysis<PostDominatorTreeWrapperPass>().getPostDomTree();
-  auto &LI = getAnalysis<LoopInfoWrapperPass>().getLoopInfo();
-  LegacyDivergenceAnalysisImpl::run(F, TTI, DT, PDT, LI);
-  LLVM_DEBUG(dbgs() << "\nAfter divergence analysis on " << F.getName()
-                    << ":\n";
-             LegacyDivergenceAnalysisImpl::print(dbgs(), F.getParent()));
-
-  return false;
-}
-
-PreservedAnalyses
-LegacyDivergenceAnalysisPass::run(Function &F, FunctionAnalysisManager &AM) {
-  auto &TTI = AM.getResult<TargetIRAnalysis>(F);
-  if (!TTI.hasBranchDivergence())
-    return PreservedAnalyses::all();
-
-  DivergentValues.clear();
-  DivergentUses.clear();
-  gpuDA = nullptr;
-
-  auto &DT = AM.getResult<DominatorTreeAnalysis>(F);
-  auto &PDT = AM.getResult<PostDominatorTreeAnalysis>(F);
-  auto &LI = AM.getResult<LoopAnalysis>(F);
-  LegacyDivergenceAnalysisImpl::run(F, TTI, DT, PDT, LI);
-  LLVM_DEBUG(dbgs() << "\nAfter divergence analysis on " << F.getName()
-                    << ":\n";
-             LegacyDivergenceAnalysisImpl::print(dbgs(), F.getParent()));
-  return PreservedAnalyses::all();
-}

llvm/lib/Analysis/SyncDependenceAnalysis.cpp

 //===--- SyncDependenceAnalysis.cpp - Compute Control Divergence Effects --===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file implements an algorithm that returns for a divergent branch
 // the set of basic blocks whose phi nodes become divergent due to divergent
 // control. These are the blocks that are reachable by two disjoint paths from
 // the branch or loop exits that have a reaching path that is disjoint from a
 // path to the loop latch.
 //
-// The SyncDependenceAnalysis is used in the DivergenceAnalysis to model
+// The SyncDependenceAnalysis is used in the UniformityAnalysis to model
 // control-induced divergence in phi nodes.
 //
 //
 // -- Reference --
-// The algorithm is presented in Section 5 of 
+// The algorithm is presented in Section 5 of
 //
 //   An abstract interpretation for SPMD divergence
 //       on reducible control flow graphs.
 //   Julian Rosemann, Simon Moll and Sebastian Hack
 //   POPL '21
 //
 //
 // -- Sync dependence --

 //     x2 = phi   E
 //         \     /
 //         x3 = phi
 //
 // The blocks D and F contain phi nodes and are thus each reachable
 // by two disjoins paths from A.
 //
 // -- Remarks --
-// * In case of loop exits we need to check the disjoint path criterion for loops.
+// * In case of loop exits we need to check the disjoint path criterion for
+// loops.
 //   To this end, we check whether the definition of x differs between the
 //   loop exit and the loop header (_after_ SSA construction).
 //
 // -- Known Limitations & Future Work --
 // * The algorithm requires reducible loops because the implementation
-//   implicitly performs a single iteration of the underlying data flow analysis.
-//   This was done for pragmatism, simplicity and speed.
+//   implicitly performs a single iteration of the underlying data flow
+//   analysis. This was done for pragmatism, simplicity and speed.
 //
 //   Relevant related work for extending the algorithm to irreducible control:
 //     A simple algorithm for global data flow analysis problems.
 //     Matthew S. Hecht and Jeffrey D. Ullman.
 //     SIAM Journal on Computing, 4(4):519–532, December 1975.
 //
 // * Another reason for requiring reducible loops is that points of
 //   synchronization in irreducible loops aren't 'obvious' - there is no unique

 using namespace llvm;
 
 // Custom Post-Order Traveral
 //
 // We cannot use the vanilla (R)PO computation of LLVM because:
 // * We (virtually) modify the CFG.
 // * We want a loop-compact block enumeration, that is the numbers assigned to
 //   blocks of a loop form an interval
-//   
+//
 using POCB = std::function<void(const BasicBlock &)>;
 using VisitedSet = std::set<const BasicBlock *>;
 using BlockStack = std::vector<const BasicBlock *>;
 
 // forward
 static void computeLoopPO(const LoopInfo &LI, Loop &Loop, POCB CallBack,
                           VisitedSet &Finalized);
 

llvm/lib/Analysis/TargetTransformInfo.cpp

return PredictableBranchThreshold.getNumOccurrences() > 0
              ? BranchProbability(PredictableBranchThreshold, 100)
              : TTIImpl->getPredictableBranchThreshold();
 }
 
 bool TargetTransformInfo::hasBranchDivergence() const {
   return TTIImpl->hasBranchDivergence();
 }
 
-bool TargetTransformInfo::useGPUDivergenceAnalysis() const {
-  return TTIImpl->useGPUDivergenceAnalysis();
-}
-
 bool TargetTransformInfo::isSourceOfDivergence(const Value *V) const {
   return TTIImpl->isSourceOfDivergence(V);
 }
 
 bool llvm::TargetTransformInfo::isAlwaysUniform(const Value *V) const {
   return TTIImpl->isAlwaysUniform(V);
 }
 

llvm/lib/Passes/PassBuilder.cpp

 #include "llvm/Analysis/CallPrinter.h"
 #include "llvm/Analysis/CostModel.h"
 #include "llvm/Analysis/CycleAnalysis.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/DomPrinter.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"
 #include "llvm/Analysis/LazyCallGraph.h"
 #include "llvm/Analysis/LazyValueInfo.h"
-#include "llvm/Analysis/LegacyDivergenceAnalysis.h"
 #include "llvm/Analysis/Lint.h"
 #include "llvm/Analysis/LoopAccessAnalysis.h"
 #include "llvm/Analysis/LoopCacheAnalysis.h"
 #include "llvm/Analysis/LoopInfo.h"
 #include "llvm/Analysis/LoopNestAnalysis.h"
 #include "llvm/Analysis/MemDerefPrinter.h"
 #include "llvm/Analysis/MemoryDependenceAnalysis.h"
 #include "llvm/Analysis/MemorySSA.h"

llvm/lib/Passes/PassRegistry.def

 FUNCTION_ANALYSIS("should-not-run-function-passes", ShouldNotRunFunctionPassesAnalysis())
 FUNCTION_ANALYSIS("should-run-extra-vector-passes", ShouldRunExtraVectorPasses())
 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())
 FUNCTION_ANALYSIS("uniformity", UniformityInfoAnalysis())
 
 #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("objc-arc-aa", objcarc::ObjCARCAA())

 FUNCTION_PASS("invalidate<all>", InvalidateAllAnalysesPass())
 FUNCTION_PASS("irce", IRCEPass())
 FUNCTION_PASS("float2int", Float2IntPass())
 FUNCTION_PASS("no-op-function", NoOpFunctionPass())
 FUNCTION_PASS("libcalls-shrinkwrap", LibCallsShrinkWrapPass())
 FUNCTION_PASS("lint", LintPass())
 FUNCTION_PASS("inject-tli-mappings", InjectTLIMappings())
 FUNCTION_PASS("instnamer", InstructionNamerPass())
-FUNCTION_PASS("legacy-divergence-analysis", LegacyDivergenceAnalysisPass())
 FUNCTION_PASS("loweratomic", LowerAtomicPass())
 FUNCTION_PASS("lower-expect", LowerExpectIntrinsicPass())
 FUNCTION_PASS("lower-guard-intrinsic", LowerGuardIntrinsicPass())
 FUNCTION_PASS("lower-constant-intrinsics", LowerConstantIntrinsicsPass())
 FUNCTION_PASS("lower-widenable-condition", LowerWidenableConditionPass())
 FUNCTION_PASS("guard-widening", GuardWideningPass())
 FUNCTION_PASS("load-store-vectorizer", LoadStoreVectorizerPass())
 FUNCTION_PASS("loop-simplify", LoopSimplifyPass())

 FUNCTION_PASS("place-safepoints", PlaceSafepointsPass())
 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<cost-model>", CostModelPrinterPass(dbgs()))
 FUNCTION_PASS("print<cycles>", CycleInfoPrinterPass(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/read_register.ll

-; RUN: opt -mtriple amdgcn-unknown-amdhsa -mcpu=gfx90a -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
-
-; CHECK-LABEL: Divergence Analysis' for function 'read_register_exec':
-; CHECK-NOT: DIVERGENT
-define i64 @read_register_exec() {
-  %reg = call i64 @llvm.read_register.i64(metadata !0)
-  ret i64 %reg
-}
-
-; CHECK-LABEL: Divergence Analysis' for function 'read_register_m0':
-; CHECK-NOT: DIVERGENT
-define i32 @read_register_m0() {
-  %reg = call i32 @llvm.read_register.i32(metadata !1)
-  ret i32 %reg
-}
-
-; CHECK-LABEL: Divergence Analysis' for function 'read_register_s17':
-; CHECK-NOT: DIVERGENT
-define i32 @read_register_s17() {
-  %reg = call i32 @llvm.read_register.i32(metadata !2)
-  ret i32 %reg
-}
-
-; CHECK-LABEL: Divergence Analysis' for function 'read_register_s17_i17':
-; CHECK-NOT: DIVERGENT
-define i17 @read_register_s17_i17() {
-  %reg = call i17 @llvm.read_register.i17(metadata !2)
-  ret i17 %reg
-}
-
-; CHECK-LABEL: Divergence Analysis' for function 'read_register_v0':
-; CHECK: DIVERGENT
-define i32 @read_register_v0() {
-  %reg = call i32 @llvm.read_register.i32(metadata !3)
-  ret i32 %reg
-}
-
-; CHECK-LABEL: Divergence Analysis' for function 'read_register_v0_v1':
-; CHECK: DIVERGENT
-define i64 @read_register_v0_v1() {
-  %reg = call i64 @llvm.read_register.i64(metadata !4)
-  ret i64 %reg
-}
-
-; CHECK-LABEL: Divergence Analysis' for function 'read_register_a0':
-; CHECK: DIVERGENT
-define i32 @read_register_a0() {
-  %reg = call i32 @llvm.read_register.i32(metadata !5)
-  ret i32 %reg
-}
-
-; CHECK-LABEL: Divergence Analysis' for function 'read_register_a0_a1':
-; CHECK: DIVERGENT
-define i64 @read_register_a0_a1() {
-  %reg = call i64 @llvm.read_register.i64(metadata !6)
-  ret i64 %reg
-}
-
-; CHECK-LABEL: Divergence Analysis' for function 'read_register_vcc_i64':
-; CHECK-NOT: DIVERGENT
-define i64 @read_register_vcc_i64() {
-  %reg = call i64 @llvm.read_register.i64(metadata !7)
-  ret i64 %reg
-}
-
-; CHECK-LABEL: Divergence Analysis' for function 'read_register_vcc_i1':
-; CHECK: DIVERGENT
-define i1 @read_register_vcc_i1() {
-  %reg = call i1 @llvm.read_register.i1(metadata !7)
-  ret i1 %reg
-}
-
-; CHECK-LABEL: Divergence Analysis' for function 'read_register_invalid_reg':
-; CHECK-NOT: DIVERGENT
-define i64 @read_register_invalid_reg() {
-  %reg = call i64 @llvm.read_register.i64(metadata !8)
-  ret i64 %reg
-}
-
-; CHECK-LABEL: Divergence Analysis' for function 'read_register_flat_scratch':
-; CHECK-NOT: DIVERGENT
-define i32 @read_register_flat_scratch() {
-  %reg = call i32 @llvm.read_register.i32(metadata !9)
-  ret i32 %reg
-}
-
-; CHECK-LABEL: Divergence Analysis' for function 'read_register_vcc_lo_i32':
-; CHECK-NOT: DIVERGENT
-define i32 @read_register_vcc_lo_i32() {
-  %reg = call i32 @llvm.read_register.i32(metadata !10)
-  ret i32 %reg
-}
-
-; CHECK-LABEL: Divergence Analysis' for function 'read_register_vcc_hi_i32':
-; CHECK-NOT: DIVERGENT
-define i32 @read_register_vcc_hi_i32() {
-  %reg = call i32 @llvm.read_register.i32(metadata !11)
-  ret i32 %reg
-}
-
-; CHECK-LABEL: Divergence Analysis' for function 'read_register_exec_lo_i32':
-; CHECK-NOT: DIVERGENT
-define i32 @read_register_exec_lo_i32() {
-  %reg = call i32 @llvm.read_register.i32(metadata !12)
-  ret i32 %reg
-}
-
-; CHECK-LABEL: Divergence Analysis' for function 'read_register_exec_hi_i32':
-; CHECK-NOT: DIVERGENT
-define i32 @read_register_exec_hi_i32() {
-  %reg = call i32 @llvm.read_register.i32(metadata !13)
-  ret i32 %reg
-}
-
-; FIXME: Why does the verifier allow this?
-; CHECK-LABEL: Divergence Analysis' for function 'read_register_empty_str_i32':
-; CHECK-NOT: DIVERGENT
-define i32 @read_register_empty_str_i32() {
-  %reg = call i32 @llvm.read_register.i32(metadata !14)
-  ret i32 %reg
-}
-
-declare i64 @llvm.read_register.i64(metadata)
-declare i32 @llvm.read_register.i32(metadata)
-declare i17 @llvm.read_register.i17(metadata)
-declare i1 @llvm.read_register.i1(metadata)
-
-!0 = !{!"exec"}
-!1 = !{!"m0"}
-!2 = !{!"s17"}
-!3 = !{!"v0"}
-!4 = !{!"v[0:1]"}
-!5 = !{!"a0"}
-!6 = !{!"a[0:1]"}
-!7 = !{!"vcc"}
-!8 = !{!"not a register"}
-!9 = !{!"flat_scratch"}
-!10 = !{!"vcc_lo"}
-!11 = !{!"vcc_hi"}
-!12 = !{!"exec_lo"}
-!13 = !{!"exec_hi"}
-!14 = !{!""}

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

-; RUN: opt -mtriple=amdgcn-- -passes='print<divergence>' 2>&1 -disable-output %s | FileCheck %s
-
-; CHECK: DIVERGENT: %orig = atomicrmw xchg ptr %ptr, i32 %val seq_cst
-define amdgpu_kernel void @test1(ptr %ptr, i32 %val) #0 {
-  %orig = atomicrmw xchg ptr %ptr, i32 %val seq_cst
-  store i32 %orig, ptr %ptr
-  ret void
-}
-
-; CHECK: DIVERGENT: %orig = cmpxchg ptr %ptr, i32 %cmp, i32 %new seq_cst seq_cst
-define amdgpu_kernel void @test2(ptr %ptr, i32 %cmp, i32 %new) {
-  %orig = cmpxchg ptr %ptr, i32 %cmp, i32 %new seq_cst seq_cst
-  %val = extractvalue { i32, i1 } %orig, 0
-  store i32 %val, ptr %ptr
-  ret void
-}
-
-; CHECK: DIVERGENT: %ret = call i32 @llvm.amdgcn.atomic.inc.i32.p1(ptr addrspace(1) %ptr, i32 %val, i32 0, i32 0, i1 false)
-define i32 @test_atomic_inc_i32(ptr addrspace(1) %ptr, i32 %val) #0 {
-  %ret = call i32 @llvm.amdgcn.atomic.inc.i32.p1(ptr addrspace(1) %ptr, i32 %val, i32 0, i32 0, i1 false)
-  ret i32 %ret
-}
-
-; CHECK: DIVERGENT: %ret = call i64 @llvm.amdgcn.atomic.inc.i64.p1(ptr addrspace(1) %ptr, i64 %val, i32 0, i32 0, i1 false)
-define i64 @test_atomic_inc_i64(ptr addrspace(1) %ptr, i64 %val) #0 {
-  %ret = call i64 @llvm.amdgcn.atomic.inc.i64.p1(ptr addrspace(1) %ptr, i64 %val, i32 0, i32 0, i1 false)
-  ret i64 %ret
-}
-
-; CHECK: DIVERGENT: %ret = call i32 @llvm.amdgcn.atomic.dec.i32.p1(ptr addrspace(1) %ptr, i32 %val, i32 0, i32 0, i1 false)
-define i32 @test_atomic_dec_i32(ptr addrspace(1) %ptr, i32 %val) #0 {
-  %ret = call i32 @llvm.amdgcn.atomic.dec.i32.p1(ptr addrspace(1) %ptr, i32 %val, i32 0, i32 0, i1 false)
-  ret i32 %ret
-}
-
-; CHECK: DIVERGENT: %ret = call i64 @llvm.amdgcn.atomic.dec.i64.p1(ptr addrspace(1) %ptr, i64 %val, i32 0, i32 0, i1 false)
-define i64 @test_atomic_dec_i64(ptr addrspace(1) %ptr, i64 %val) #0 {
-  %ret = call i64 @llvm.amdgcn.atomic.dec.i64.p1(ptr addrspace(1) %ptr, i64 %val, i32 0, i32 0, i1 false)
-  ret i64 %ret
-}
-
-declare i32 @llvm.amdgcn.atomic.inc.i32.p1(ptr addrspace(1) nocapture, i32, i32, i32, i1) #1
-declare i64 @llvm.amdgcn.atomic.inc.i64.p1(ptr addrspace(1) nocapture, i64, i32, i32, i1) #1
-declare i32 @llvm.amdgcn.atomic.dec.i32.p1(ptr addrspace(1) nocapture, i32, i32, i32, i1) #1
-declare i64 @llvm.amdgcn.atomic.dec.i64.p1(ptr addrspace(1) nocapture, i64, i32, i32, i1) #1
-
-; CHECK: DIVERGENT: %ret = call i32 @llvm.amdgcn.global.atomic.csub.p1(ptr addrspace(1) %ptr, i32 %val)
-define amdgpu_kernel void @test_atomic_csub_i32(ptr addrspace(1) %ptr, i32 %val) #0 {
-  %ret = call i32 @llvm.amdgcn.global.atomic.csub.p1(ptr addrspace(1) %ptr, i32 %val)
-  store i32 %ret, ptr addrspace(1) %ptr, align 4
-  ret void
-}
-
-declare i32 @llvm.amdgcn.global.atomic.csub.p1(ptr addrspace(1) nocapture, i32) #1
-
-attributes #0 = { nounwind }
-attributes #1 = { argmemonly nounwind willreturn }

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

-; RUN: opt -mtriple=amdgcn-- -passes='print<divergence>' 2>&1 -disable-output %s | FileCheck %s
-
-; CHECK: DIVERGENT: %swizzle = call i32 @llvm.amdgcn.ds.swizzle(i32 %src, i32 100) #0
-define amdgpu_kernel void @ds_swizzle(ptr addrspace(1) %out, i32 %src) #0 {
-  %swizzle = call i32 @llvm.amdgcn.ds.swizzle(i32 %src, i32 100) #0
-  store i32 %swizzle, ptr 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-- -passes='print<divergence>' 2>&1 -disable-output | FileCheck %s
-
-; CHECK-LABEL: function 'test_amdgpu_ps':
-; CHECK: DIVERGENT:  ptr 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(ptr 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: 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(ptr 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: function 'test_c':
-; CHECK: DIVERGENT:
-; CHECK: DIVERGENT:
-; CHECK: DIVERGENT:
-; CHECK: DIVERGENT:
-; CHECK: DIVERGENT:
-; CHECK: DIVERGENT:
-; CHECK: DIVERGENT:
-define void @test_c(ptr addrspace(4) 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/LegacyDivergenceAnalysis/AMDGPU/lit.local.cfg

-if not 'AMDGPU' in config.root.targets:
-    config.unsupported = True

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

-;RUN: opt -mtriple=amdgcn-mesa-mesa3d -passes='print<divergence>' 2>&1 -disable-output %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
-}
-
-;CHECK: DIVERGENT: %orig = call i32 @llvm.amdgcn.buffer.atomic.add.i32(
-define float @buffer_atomic_add(<4 x i32> inreg %rsrc, i32 inreg %data) #0 {
-main_body:
-  %orig = call i32 @llvm.amdgcn.buffer.atomic.add.i32(i32 %data, <4 x i32> %rsrc, i32 0, i32 0, i1 0)
-  %r = bitcast i32 %orig to float
-  ret float %r
-}
-
-;CHECK: DIVERGENT: %orig = call i32 @llvm.amdgcn.buffer.atomic.sub.i32(
-define float @buffer_atomic_sub(<4 x i32> inreg %rsrc, i32 inreg %data) #0 {
-main_body:
-  %orig = call i32 @llvm.amdgcn.buffer.atomic.sub.i32(i32 %data, <4 x i32> %rsrc, i32 0, i32 0, i1 0)
-  %r = bitcast i32 %orig to float
-  ret float %r
-}
-
-;CHECK: DIVERGENT: %orig = call i32 @llvm.amdgcn.buffer.atomic.smin.i32(
-define float @buffer_atomic_smin(<4 x i32> inreg %rsrc, i32 inreg %data) #0 {
-main_body:
-  %orig = call i32 @llvm.amdgcn.buffer.atomic.smin.i32(i32 %data, <4 x i32> %rsrc, i32 0, i32 0, i1 0)
-  %r = bitcast i32 %orig to float
-  ret float %r
-}
-
-;CHECK: DIVERGENT: %orig = call i32 @llvm.amdgcn.buffer.atomic.umin.i32(
-define float @buffer_atomic_umin(<4 x i32> inreg %rsrc, i32 inreg %data) #0 {
-main_body:
-  %orig = call i32 @llvm.amdgcn.buffer.atomic.umin.i32(i32 %data, <4 x i32> %rsrc, i32 0, i32 0, i1 0)
-  %r = bitcast i32 %orig to float
-  ret float %r
-}
-
-;CHECK: DIVERGENT: %orig = call i32 @llvm.amdgcn.buffer.atomic.smax.i32(
-define float @buffer_atomic_smax(<4 x i32> inreg %rsrc, i32 inreg %data) #0 {
-main_body:
-  %orig = call i32 @llvm.amdgcn.buffer.atomic.smax.i32(i32 %data, <4 x i32> %rsrc, i32 0, i32 0, i1 0)
-  %r = bitcast i32 %orig to float
-  ret float %r
-}
-
-;CHECK: DIVERGENT: %orig = call i32 @llvm.amdgcn.buffer.atomic.umax.i32(
-define float @buffer_atomic_umax(<4 x i32> inreg %rsrc, i32 inreg %data) #0 {
-main_body:
-  %orig = call i32 @llvm.amdgcn.buffer.atomic.umax.i32(i32 %data, <4 x i32> %rsrc, i32 0, i32 0, i1 0)
-  %r = bitcast i32 %orig to float
-  ret float %r
-}
-
-;CHECK: DIVERGENT: %orig = call i32 @llvm.amdgcn.buffer.atomic.and.i32(
-define float @buffer_atomic_and(<4 x i32> inreg %rsrc, i32 inreg %data) #0 {
-main_body:
-  %orig = call i32 @llvm.amdgcn.buffer.atomic.and.i32(i32 %data, <4 x i32> %rsrc, i32 0, i32 0, i1 0)
-  %r = bitcast i32 %orig to float
-  ret float %r
-}
-
-;CHECK: DIVERGENT: %orig = call i32 @llvm.amdgcn.buffer.atomic.or.i32(
-define float @buffer_atomic_or(<4 x i32> inreg %rsrc, i32 inreg %data) #0 {
-main_body:
-  %orig = call i32 @llvm.amdgcn.buffer.atomic.or.i32(i32 %data, <4 x i32> %rsrc, i32 0, i32 0, i1 0)
-  %r = bitcast i32 %orig to float
-  ret float %r
-}
-
-;CHECK: DIVERGENT: %orig = call i32 @llvm.amdgcn.buffer.atomic.xor.i32(
-define float @buffer_atomic_xor(<4 x i32> inreg %rsrc, i32 inreg %data) #0 {
-main_body:
-  %orig = call i32 @llvm.amdgcn.buffer.atomic.xor.i32(i32 %data, <4 x i32> %rsrc, i32 0, i32 0, i1 0)
-  %r = bitcast i32 %orig to float
-  ret float %r
-}
-
-;CHECK: DIVERGENT: %orig = call i32 @llvm.amdgcn.buffer.atomic.cmpswap(
-define float @buffer_atomic_cmpswap(<4 x i32> inreg %rsrc, i32 inreg %data, i32 inreg %cmp) #0 {
-main_body:
-  %orig = call i32 @llvm.amdgcn.buffer.atomic.cmpswap(i32 %data, i32 %cmp, <4 x i32> %rsrc, i32 0, i32 0, i1 0)
-  %r = bitcast i32 %orig to float
-  ret float %r
-}
-
-declare i32 @llvm.amdgcn.buffer.atomic.swap.i32(i32, <4 x i32>, i32, i32, i1) #0
-declare i32 @llvm.amdgcn.buffer.atomic.add.i32(i32, <4 x i32>, i32, i32, i1) #0
-declare i32 @llvm.amdgcn.buffer.atomic.sub.i32(i32, <4 x i32>, i32, i32, i1) #0
-declare i32 @llvm.amdgcn.buffer.atomic.smin.i32(i32, <4 x i32>, i32, i32, i1) #0
-declare i32 @llvm.amdgcn.buffer.atomic.umin.i32(i32, <4 x i32>, i32, i32, i1) #0
-declare i32 @llvm.amdgcn.buffer.atomic.smax.i32(i32, <4 x i32>, i32, i32, i1) #0
-declare i32 @llvm.amdgcn.buffer.atomic.umax.i32(i32, <4 x i32>, i32, i32, i1) #0
-declare i32 @llvm.amdgcn.buffer.atomic.and.i32(i32, <4 x i32>, i32, i32, i1) #0
-declare i32 @llvm.amdgcn.buffer.atomic.or.i32(i32, <4 x i32>, i32, i32, i1) #0
-declare i32 @llvm.amdgcn.buffer.atomic.xor.i32(i32, <4 x i32>, i32, i32, i1) #0
-declare i32 @llvm.amdgcn.buffer.atomic.cmpswap(i32, i32, <4 x i32>, i32, i32, i1) #0
-
-attributes #0 = { nounwind }

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

-;RUN: opt -mtriple=amdgcn-mesa-mesa3d -passes='print<divergence>' 2>&1 -disable-output %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
-}
-
-;CHECK: DIVERGENT: %orig = call i32 @llvm.amdgcn.image.atomic.add.1d.i32.i32(
-define float @image_atomic_add(<8 x i32> inreg %rsrc, i32 inreg %addr, i32 inreg %data) #0 {
-main_body:
-  %orig = call i32 @llvm.amdgcn.image.atomic.add.1d.i32.i32(i32 %data, i32 %addr, <8 x i32> %rsrc, i32 0, i32 0)
-  %r = bitcast i32 %orig to float
-  ret float %r
-}
-
-;CHECK: DIVERGENT: %orig = call i32 @llvm.amdgcn.image.atomic.sub.1d.i32.i32(
-define float @image_atomic_sub(<8 x i32> inreg %rsrc, i32 inreg %addr, i32 inreg %data) #0 {
-main_body:
-  %orig = call i32 @llvm.amdgcn.image.atomic.sub.1d.i32.i32(i32 %data, i32 %addr, <8 x i32> %rsrc, i32 0, i32 0)
-  %r = bitcast i32 %orig to float
-  ret float %r
-}
-
-;CHECK: DIVERGENT: %orig = call i32 @llvm.amdgcn.image.atomic.smin.1d.i32.i32(
-define float @image_atomic_smin(<8 x i32> inreg %rsrc, i32 inreg %addr, i32 inreg %data) #0 {
-main_body:
-  %orig = call i32 @llvm.amdgcn.image.atomic.smin.1d.i32.i32(i32 %data, i32 %addr, <8 x i32> %rsrc, i32 0, i32 0)
-  %r = bitcast i32 %orig to float
-  ret float %r
-}
-
-;CHECK: DIVERGENT: %orig = call i32 @llvm.amdgcn.image.atomic.umin.1d.i32.i32(
-define float @image_atomic_umin(<8 x i32> inreg %rsrc, i32 inreg %addr, i32 inreg %data) #0 {
-main_body:
-  %orig = call i32 @llvm.amdgcn.image.atomic.umin.1d.i32.i32(i32 %data, i32 %addr, <8 x i32> %rsrc, i32 0, i32 0)
-  %r = bitcast i32 %orig to float
-  ret float %r
-}
-
-;CHECK: DIVERGENT: %orig = call i32 @llvm.amdgcn.image.atomic.smax.1d.i32.i32(
-define float @image_atomic_smax(<8 x i32> inreg %rsrc, i32 inreg %addr, i32 inreg %data) #0 {
-main_body:
-  %orig = call i32 @llvm.amdgcn.image.atomic.smax.1d.i32.i32(i32 %data, i32 %addr, <8 x i32> %rsrc, i32 0, i32 0)
-  %r = bitcast i32 %orig to float
-  ret float %r
-}
-
-;CHECK: DIVERGENT: %orig = call i32 @llvm.amdgcn.image.atomic.umax.1d.i32.i32(
-define float @image_atomic_umax(<8 x i32> inreg %rsrc, i32 inreg %addr, i32 inreg %data) #0 {
-main_body:
-  %orig = call i32 @llvm.amdgcn.image.atomic.umax.1d.i32.i32(i32 %data, i32 %addr, <8 x i32> %rsrc, i32 0, i32 0)
-  %r = bitcast i32 %orig to float
-  ret float %r
-}
-
-;CHECK: DIVERGENT: %orig = call i32 @llvm.amdgcn.image.atomic.and.1d.i32.i32(
-define float @image_atomic_and(<8 x i32> inreg %rsrc, i32 inreg %addr, i32 inreg %data) #0 {
-main_body:
-  %orig = call i32 @llvm.amdgcn.image.atomic.and.1d.i32.i32(i32 %data, i32 %addr, <8 x i32> %rsrc, i32 0, i32 0)
-  %r = bitcast i32 %orig to float
-  ret float %r
-}
-
-;CHECK: DIVERGENT: %orig = call i32 @llvm.amdgcn.image.atomic.or.1d.i32.i32(
-define float @image_atomic_or(<8 x i32> inreg %rsrc, i32 inreg %addr, i32 inreg %data) #0 {
-main_body:
-  %orig = call i32 @llvm.amdgcn.image.atomic.or.1d.i32.i32(i32 %data, i32 %addr, <8 x i32> %rsrc, i32 0, i32 0)
-  %r = bitcast i32 %orig to float
-  ret float %r
-}
-
-;CHECK: DIVERGENT: %orig = call i32 @llvm.amdgcn.image.atomic.xor.1d.i32.i32(
-define float @image_atomic_xor(<8 x i32> inreg %rsrc, i32 inreg %addr, i32 inreg %data) #0 {
-main_body:
-  %orig = call i32 @llvm.amdgcn.image.atomic.xor.1d.i32.i32(i32 %data, i32 %addr, <8 x i32> %rsrc, i32 0, i32 0)
-  %r = bitcast i32 %orig to float
-  ret float %r
-}
-
-;CHECK: DIVERGENT: %orig = call i32 @llvm.amdgcn.image.atomic.inc.1d.i32.i32(
-define float @image_atomic_inc(<8 x i32> inreg %rsrc, i32 inreg %addr, i32 inreg %data) #0 {
-main_body:
-  %orig = call i32 @llvm.amdgcn.image.atomic.inc.1d.i32.i32(i32 %data, i32 %addr, <8 x i32> %rsrc, i32 0, i32 0)
-  %r = bitcast i32 %orig to float
-  ret float %r
-}
-
-;CHECK: DIVERGENT: %orig = call i32 @llvm.amdgcn.image.atomic.dec.1d.i32.i32(
-define float @image_atomic_dec(<8 x i32> inreg %rsrc, i32 inreg %addr, i32 inreg %data) #0 {
-main_body:
-  %orig = call i32 @llvm.amdgcn.image.atomic.dec.1d.i32.i32(i32 %data, i32 %addr, <8 x i32> %rsrc, i32 0, i32 0)
-  %r = bitcast i32 %orig to float
-  ret float %r
-}
-
-;CHECK: DIVERGENT: %orig = call i32 @llvm.amdgcn.image.atomic.cmpswap.1d.i32.i32(
-define float @image_atomic_cmpswap(<8 x i32> inreg %rsrc, i32 inreg %addr, i32 inreg %data, i32 inreg %cmp) #0 {
-main_body:
-  %orig = call i32 @llvm.amdgcn.image.atomic.cmpswap.1d.i32.i32(i32 %data, i32 %cmp, i32 %addr, <8 x i32> %rsrc, i32 0, i32 0)
-  %r = bitcast i32 %orig to float
-  ret float %r
-}
-
-;CHECK: DIVERGENT: %orig = call i32 @llvm.amdgcn.image.atomic.add.2d.i32.i32(
-define float @image_atomic_add_2d(<8 x i32> inreg %rsrc, i32 inreg %s, i32 inreg %t, i32 inreg %data) #0 {
-main_body:
-  %orig = call i32 @llvm.amdgcn.image.atomic.add.2d.i32.i32(i32 %data, i32 %s, i32 %t, <8 x i32> %rsrc, i32 0, i32 0)
-  %r = bitcast i32 %orig to float
-  ret float %r
-}
-
-declare i32 @llvm.amdgcn.image.atomic.swap.1d.i32.i32(i32, i32, <8 x i32>, i32, i32) #0
-declare i32 @llvm.amdgcn.image.atomic.add.1d.i32.i32(i32, i32, <8 x i32>, i32, i32) #0
-declare i32 @llvm.amdgcn.image.atomic.sub.1d.i32.i32(i32, i32, <8 x i32>, i32, i32) #0
-declare i32 @llvm.amdgcn.image.atomic.smin.1d.i32.i32(i32, i32, <8 x i32>, i32, i32) #0
-declare i32 @llvm.amdgcn.image.atomic.umin.1d.i32.i32(i32, i32, <8 x i32>, i32, i32) #0
-declare i32 @llvm.amdgcn.image.atomic.smax.1d.i32.i32(i32, i32, <8 x i32>, i32, i32) #0
-declare i32 @llvm.amdgcn.image.atomic.umax.1d.i32.i32(i32, i32, <8 x i32>, i32, i32) #0
-declare i32 @llvm.amdgcn.image.atomic.and.1d.i32.i32(i32, i32, <8 x i32>, i32, i32) #0
-declare i32 @llvm.amdgcn.image.atomic.or.1d.i32.i32(i32, i32, <8 x i32>, i32, i32) #0
-declare i32 @llvm.amdgcn.image.atomic.xor.1d.i32.i32(i32, i32, <8 x i32>, i32, i32) #0
-declare i32 @llvm.amdgcn.image.atomic.inc.1d.i32.i32(i32, i32, <8 x i32>, i32, i32) #0
-declare i32 @llvm.amdgcn.image.atomic.dec.1d.i32.i32(i32, i32, <8 x i32>, i32, i32) #0
-declare i32 @llvm.amdgcn.image.atomic.cmpswap.1d.i32.i32(i32, i32, i32, <8 x i32>, i32, i32) #0
-
-declare i32 @llvm.amdgcn.image.atomic.add.2d.i32.i32(i32, i32, i32, <8 x i32>, i32, i32) #0
-
-attributes #0 = { nounwind }

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

-; RUN: opt -mtriple=amdgcn-- -passes='print<divergence>' 2>&1 -disable-output %s | FileCheck %s
-
-; Test that we consider loads from flat and private addrspaces to be divergent.
-
-; CHECK: DIVERGENT: %val = load i32, ptr %flat, align 4
-define amdgpu_kernel void @flat_load(ptr %flat) {
-  %val = load i32, ptr %flat, align 4
-  ret void
-}
-
-; CHECK: DIVERGENT: %val = load i32, ptr addrspace(5) %priv, align 4
-define amdgpu_kernel void @private_load(ptr addrspace(5) %priv) {
-  %val = load i32, ptr addrspace(5) %priv, align 4
-  ret void
-}

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

-; RUN: opt %s -mtriple amdgcn-- -passes='print<divergence>' 2>&1 -disable-output | FileCheck %s
-
-; CHECK: DIVERGENT:  %tmp5 = getelementptr inbounds float, ptr addrspace(1) %arg, i64 %tmp2
-; CHECK: DIVERGENT:  %tmp10 = load volatile float, ptr addrspace(1) %tmp5, align 4
-; CHECK: DIVERGENT:  %tmp11 = load volatile float, ptr 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(ptr addrspace(1) noalias nocapture readonly %arg, ptr addrspace(1) noalias nocapture readonly %arg1) #0 {
-bb0:
-  %tmp = tail call i32 @llvm.amdgcn.workitem.id.x() #0
-  %tmp2 = sext i32 %tmp to i64
-  %tmp5 = getelementptr inbounds float, ptr addrspace(1) %arg, i64 %tmp2
-  %tmp6 = load volatile float, ptr addrspace(1) %tmp5, align 4
-  %tmp8 = fcmp olt float %tmp6, 0.000000e+00
-  br i1 %tmp8, label %bb1, label %bb2
-
-bb1:
-  %tmp10 = load volatile float, ptr addrspace(1) %tmp5, align 4
-  br label %bb2
-
-bb2:
-  %tmp11 = load volatile float, ptr addrspace(1) %tmp5, align 4
-  br label %bb1
-}
-
-; Function Attrs: nounwind readnone
-declare i32 @llvm.amdgcn.workitem.id.x() #1
-
-attributes #0 = { nounwind }
-attributes #1 = { nounwind readnone }

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

-; RUN: opt -mtriple=amdgcn-- -passes='print<divergence>' 2>&1 -disable-output %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 =
-; CHECK: DIVERGENT: br i1 %break, label %end, label %header
-; Note: %counter is not divergent!
-define amdgpu_ps void @test1(i32 %bound) {
-entry:
-  br label %header
-
-header:
-  %counter = phi i32 [ 0, %entry ], [ %counter.footer, %footer ]
-  %break = icmp sge i32 %counter, %bound
-  br i1 %break, label %footer, label %body
-
-body:
-  %counter.next = add i32 %counter, 1
-  br label %footer
-
-footer:
-  %counter.footer = phi i32 [ %counter.next, %body ], [ undef, %header ]
-  br i1 %break, label %end, label %header
-
-end:
-  ret void
-}

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

-; RUN: opt %s -mtriple amdgcn-- -passes='print<divergence>' 2>&1 -disable-output | 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 ptr 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 -passes='print<divergence>' 2>&1 -disable-output %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()
-define amdgpu_kernel void @workitem_id_x() #1 {
-  %id.x = call i32 @llvm.amdgcn.workitem.id.x()
-  store volatile i32 %id.x, ptr addrspace(1) undef
-  ret void
-}
-
-; CHECK: DIVERGENT:  %id.y = call i32 @llvm.amdgcn.workitem.id.y()
-define amdgpu_kernel void @workitem_id_y() #1 {
-  %id.y = call i32 @llvm.amdgcn.workitem.id.y()
-  store volatile i32 %id.y, ptr addrspace(1) undef
-  ret void
-}
-
-; CHECK: DIVERGENT:  %id.z = call i32 @llvm.amdgcn.workitem.id.z()
-define amdgpu_kernel void @workitem_id_z() #1 {
-  %id.z = call i32 @llvm.amdgcn.workitem.id.z()
-  store volatile i32 %id.z, ptr addrspace(1) undef
-  ret void
-}
-
-; CHECK: DIVERGENT:  %mbcnt.lo = call i32 @llvm.amdgcn.mbcnt.lo(i32 0, i32 0)
-define amdgpu_kernel void @mbcnt_lo() #1 {
-  %mbcnt.lo = call i32 @llvm.amdgcn.mbcnt.lo(i32 0, i32 0)
-  store volatile i32 %mbcnt.lo, ptr addrspace(1) undef
-  ret void
-}
-
-; CHECK: DIVERGENT:  %mbcnt.hi = call i32 @llvm.amdgcn.mbcnt.hi(i32 0, i32 0)
-define amdgpu_kernel void @mbcnt_hi() #1 {
-  %mbcnt.hi = call i32 @llvm.amdgcn.mbcnt.hi(i32 0, i32 0)
-  store volatile i32 %mbcnt.hi, ptr addrspace(1) undef
-  ret void
-}
-
-attributes #0 = { nounwind readnone }
-attributes #1 = { nounwind }

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

-; RUN: opt %s -passes='print<divergence>' 2>&1 -disable-output | 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: 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: 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: 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: 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: 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: 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
-  br i1 %exit_cond, label %loop_exit, label %loop
-loop_exit:
-  %i.lcssa = phi i32 [ %i, %loop ]
-; CHECK: DIVERGENT: %i.lcssa =
-  %cond = icmp eq i32 %i.lcssa, 10
-  br i1 %cond, label %then, label %else
-; CHECK: DIVERGENT: br i1 %cond,
-then:
-  ret i32 0
-else:
-  ret i32 1
-}
-
-; 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: 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:
-  %i2 = phi i32 [ 0, %entry ], [ %i3, %loop_latch ]
-  %j2 = add i32 %i2, 2
-  br label %loop_body
-loop_body:
-  %i3 = phi i32 [ %j1, %loop_entry_1 ], [ %j2, %loop_entry_2 ]
-  br label %loop_latch
-loop_latch:
-  %div = sdiv i32 %tid, %i3
-  switch i32 %div, label %branch [ i32 1, label %loop_entry_1
-                                   i32 2, label %loop_entry_2 ]
-branch:
-  %cmp = icmp eq i32 %i3, 5
-  br i1 %cmp, label %then, label %else
-; CHECK: DIVERGENT: br i1 %cmp,
-then:
-  ret i32 0
-else:
-  ret i32 1
-}
-
-; Verifies sync-dependence is computed correctly in the absense of loops.
-define i32 @sync_no_loop(i32 %arg) {
-entry:
-  %0 = add i32 %arg, 1
-  %tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
-  %1 = icmp sge i32 %tid, 10
-  br i1 %1, label %bb1, label %bb2
-
-bb1:
-  br label %bb3
-
-bb2:
-  br label %bb3
-
-bb3:
-  %2 = add i32 %0, 2
-  ; CHECK-NOT: DIVERGENT: %2
-  ret i32 %2
-}
-
-declare i32 @llvm.nvvm.read.ptx.sreg.tid.x()
-declare i32 @llvm.nvvm.read.ptx.sreg.tid.y()
-declare i32 @llvm.nvvm.read.ptx.sreg.tid.z()
-declare i32 @llvm.nvvm.read.ptx.sreg.laneid()
-
-!nvvm.annotations = !{!0, !1, !2, !3, !4, !5}
-!0 = !{ptr @no_diverge, !"kernel", i32 1}
-!1 = !{ptr @sync, !"kernel", i32 1}
-!2 = !{ptr @mixed, !"kernel", i32 1}
-!3 = !{ptr @loop, !"kernel", i32 1}
-!4 = !{ptr @unstructured_loop, !"kernel", i32 1}
-!5 = !{ptr @sync_no_loop, !"kernel", i32 1}

llvm/test/Analysis/LegacyDivergenceAnalysis/NVPTX/lit.local.cfg

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

llvm/test/Analysis/UniformityAnalysis/AMDGPU/MIR/always-uniform-gmir.mir

-# NOTE: This file is Generic MIR translation of test/Analysis/DivergenceAnalysis/AMDGPU/always_uniform.ll test file
+# NOTE: This file is Generic MIR translation of test/Analysis/UniformityAnalysis/AMDGPU/always_uniform.ll test file
 # RUN: llc -mtriple=amdgcn-- -run-pass=print-machine-uniformity -o - %s 2>&1 | FileCheck %s
 ---
 name:            readfirstlane
 body:             |
   bb.1:
     ; CHECK-LABEL: MachineUniformityInfo for function: readfirstlane
     ; CHECK: DIVERGENT: %{{[0-9]+}}
     ; CHECK-SAME:llvm.amdgcn.workitem.id.x

llvm/test/Analysis/UniformityAnalysis/AMDGPU/MIR/always-uniform.mir

 # RUN: llc -mtriple=amdgcn-- -mcpu=gfx900 -run-pass=print-machine-uniformity -o - %s 2>&1 | FileCheck %s
 
 # readlane, readfirstlane is always uniform
 
 ---
 name:            readlane
 machineFunctionInfo:
   isEntryFunction: true
 body:             |
   bb.0:
     ; CHECK-LABEL: MachineUniformityInfo for function: readlane
     ; CHECK-NEXT: ALL VALUES UNIFORM
     %0:vgpr_32 = IMPLICIT_DEF
     %1:vgpr_32 = IMPLICIT_DEF
     %2:sgpr_32 = V_READFIRSTLANE_B32 %0, implicit $exec
     %3:sgpr_32 = V_READLANE_B32 %1, 0, implicit $exec
     $sgpr0 = V_READFIRSTLANE_B32 $vgpr0, implicit $exec
     $sgpr1 = V_READLANE_B32 $vgpr1, $sgpr0, implicit $exec
     S_ENDPGM 0
 ...
 
 # Readlane with physical register as operand
 ---
 name:            readlane2
 machineFunctionInfo:
   isEntryFunction: true
 body:             |
   bb.0:
     ; CHECK-LABEL: MachineUniformityInfo for function: readlane2
     ; CHECK-NEXT: ALL VALUES UNIFORM
     %0:vgpr_32 = IMPLICIT_DEF
     %1:vgpr_32 = IMPLICIT_DEF
     %4:sgpr_32 = V_READLANE_B32 $vgpr0, 0, implicit $exec
     $sgpr0 = V_READFIRSTLANE_B32 $vgpr0, implicit $exec
     $sgpr1 = V_READLANE_B32 $vgpr1, $sgpr0, implicit $exec
     %5:sgpr_32 = V_READFIRSTLANE_B32 $vgpr1, implicit $exec
     S_ENDPGM 0
 ...
 
 
 
 # for copy operand src = sgpr -> uniform
 ---
 name:            sgprcopy
 tracksRegLiveness: true
 machineFunctionInfo:
   isEntryFunction: true
 body:             |
   bb.0:
     ; CHECK-LABEL: MachineUniformityInfo for function: sgprcopy
     ; CHECK-NEXT: ALL VALUES UNIFORM
     liveins: $sgpr0,$sgpr1,$vgpr0
     %0:sgpr_32 = COPY $sgpr0
     %1:vgpr_32 = COPY $sgpr1
     S_ENDPGM 0
 ...

llvm/test/Analysis/UniformityAnalysis/AMDGPU/MIR/atomics-gmir.mir

 # RUN: llc -mtriple=amdgcn-- -run-pass=print-machine-uniformity -o - %s 2>&1 | FileCheck %s
 
 ---
 name:            test1
 tracksRegLiveness: true
 body:             |
   bb.1:
     %2:_(s32) = IMPLICIT_DEF
     %3:_(s32) = IMPLICIT_DEF
     %0:_(p0) = G_MERGE_VALUES %2(s32), %3(s32)
     %1:_(s32) = IMPLICIT_DEF
 
     ; CHECK: DIVERGENT
     ; CHECK-SAME: G_ATOMICRMW_XCHG
     %4:_(s32) = G_ATOMICRMW_XCHG %0(p0), %1 :: (load store seq_cst (s32))
 
     ; CHECK: DIVERGENT
     ; CHECK-SAME: G_ATOMIC_CMPXCHG_WITH_SUCCESS
     %5:_(s32), %6:_(s1) = G_ATOMIC_CMPXCHG_WITH_SUCCESS %0(p0), %1, %2 :: (load store seq_cst seq_cst (s32) )
     $vgpr0 = COPY %4(s32)
     SI_RETURN implicit $vgpr0
 ...
 
 ---
 name:            test_atomic_inc_dec
 tracksRegLiveness: true
 body:             |
   bb.1:
 
     %2:_(s32) = IMPLICIT_DEF
     %3:_(s32) = IMPLICIT_DEF
     %0:_(p1) = G_MERGE_VALUES %2(s32), %3(s32)
     %1:_(s32) = IMPLICIT_DEF
     %5:_(s64) = IMPLICIT_DEF
 
     ; CHECK: DIVERGENT: %{{[0-9]}}: %{{[0-9]}}:_(s32) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.amdgcn.atomic.inc)
     %4:_(s32) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.amdgcn.atomic.inc), %0(p1), %1(s32), 0, 0, 0 :: (load store (s32) )
 
     ; CHECK: DIVERGENT: %{{[0-9]}}: %{{[0-9]}}:_(s64) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.amdgcn.atomic.inc)
     %6:_(s64) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.amdgcn.atomic.inc), %0(p1), %5(s64), 0, 0, 0 :: (load store (s64) )
 
     ; CHECK: DIVERGENT: %{{[0-9]}}: %{{[0-9]}}:_(s32) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.amdgcn.atomic.dec)
     %7:_(s32) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.amdgcn.atomic.dec), %0(p1), %1(s32), 0, 0, 0 :: (load store (s32) )
 
     ; CHECK: DIVERGENT: %{{[0-9]}}: %{{[0-9]}}:_(s64) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.amdgcn.atomic.dec)
     %8:_(s64) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.amdgcn.atomic.dec), %0(p1), %5(s64), 0, 0, 0 :: (load store (s64) )
     $vgpr0 = COPY %4(s32)
     SI_RETURN implicit $vgpr0
 
 ...
 
 ---
 name:            test_atomics
 tracksRegLiveness: true
 body:             |
   bb.1:
 
     %2:_(s32) = IMPLICIT_DEF
     %3:_(s32) = IMPLICIT_DEF
     %0:_(p1) = G_MERGE_VALUES %2(s32), %3(s32)
     %1:_(s32) = IMPLICIT_DEF
     %5:_(s32) = IMPLICIT_DEF
 
     ; CHECK: DIVERGENT: %{{[0-9]}}: %{{[0-9]}}:_(s32) = G_ATOMICRMW_ADD
     %4:_(s32) = G_ATOMICRMW_ADD %2, %3
 
     ; CHECK: DIVERGENT: %{{[0-9]}}: %{{[0-9]}}:_(s32) = G_ATOMICRMW_SUB
     %6:_(s32) = G_ATOMICRMW_SUB %1, %5
 
     ; CHECK: DIVERGENT: %{{[0-9]}}: %{{[0-9]}}:_(s32) = G_ATOMICRMW_AND
     %7:_(s32) = G_ATOMICRMW_AND %2, %3
 
     ; CHECK: DIVERGENT: %{{[0-9]}}: %{{[0-9]}}:_(s32) = G_ATOMICRMW_NAND
     %8:_(s32) = G_ATOMICRMW_NAND %1, %5
 
     ; CHECK: DIVERGENT: %{{[0-9]}}: %{{[0-9]}}:_(s32) = G_ATOMICRMW_OR
     %9:_(s32) = G_ATOMICRMW_OR %2, %3
 
     ; CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s32) = G_ATOMICRMW_XOR
     %10:_(s32) = G_ATOMICRMW_XOR %1, %5
 
     ; CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s32) = G_ATOMICRMW_MAX
     %11:_(s32) = G_ATOMICRMW_MAX %2, %3
 
     ; CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s32) = G_ATOMICRMW_MIN
     %12:_(s32) = G_ATOMICRMW_MIN %1, %5
 
     ; CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s32) = G_ATOMICRMW_UMAX
     %13:_(s32) = G_ATOMICRMW_UMAX %2, %3
 
     ; CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s32) = G_ATOMICRMW_UMIN
     %14:_(s32) = G_ATOMICRMW_UMIN %1, %5
 
     ; CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s32) = G_ATOMICRMW_FADD
     %15:_(s32) = G_ATOMICRMW_FADD %2, %3
 
     ; CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s32) = G_ATOMICRMW_FSUB
     %16:_(s32) = G_ATOMICRMW_FSUB %1, %5
 
     ; CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s32) = G_ATOMICRMW_FMAX
     %17:_(s32) = G_ATOMICRMW_FMAX %2, %3
 
     ; CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s32) = G_ATOMICRMW_FMIN
     %18:_(s32) = G_ATOMICRMW_FMIN %1, %5
 
     $vgpr0 = COPY %4(s32)
     SI_RETURN implicit $vgpr0
 
 ...

llvm/test/Analysis/UniformityAnalysis/AMDGPU/MIR/atomics.mir

 # RUN: llc -mtriple=amdgcn-- -run-pass=print-machine-uniformity -o - %s 2>&1 | FileCheck %s
 
 ---
 name:            test1
 tracksRegLiveness: true
 body:             |
   bb.0:
     ; CHECK-LABEL: MachineUniformityInfo for function: test1
     %2:vgpr_32 = IMPLICIT_DEF
     %1:vgpr_32 = IMPLICIT_DEF
     %0:vgpr_32 = IMPLICIT_DEF
     %3:sreg_64 = REG_SEQUENCE %0, %subreg.sub0, %1, %subreg.sub1
     %5:vreg_64 = COPY %3
     %6:vreg_64 = COPY %3
     ; CHECK: DIVERGENT
     ; CHECK-SAME: FLAT_ATOMIC_SWAP_RTN
     %4:vgpr_32 = FLAT_ATOMIC_SWAP_RTN killed %5, %2, 0, 1, implicit $exec, implicit $flat_scr :: (load store seq_cst (s32))
     ; CHECK: DIVERGENT
     ; CHECK-SAME: FLAT_ATOMIC_SWAP_RTN
     %7:vgpr_32 = FLAT_ATOMIC_SWAP_RTN killed %6, %2, 0, 1, implicit $exec, implicit $flat_scr ; No memopernads
     $vgpr0 = COPY %4
     SI_RETURN implicit $vgpr0
 ...
 
 ---
 name:            test2
 tracksRegLiveness: true
 body:             |
   bb.0:
     ; CHECK-LABEL: MachineUniformityInfo for function: test2
     %3:vgpr_32 = IMPLICIT_DEF
     %2:vgpr_32 = IMPLICIT_DEF
     %1:vgpr_32 = IMPLICIT_DEF
     %0:vgpr_32 = IMPLICIT_DEF
     %4:sreg_64 = REG_SEQUENCE %0, %subreg.sub0, %1, %subreg.sub1
     %5:sreg_64 = REG_SEQUENCE %3, %subreg.sub0, %2, %subreg.sub1
     %7:vreg_64 = COPY %4
     %8:vreg_64 = COPY %5
     ; CHECK: DIVERGENT
     ; CHECK-SAME: FLAT_ATOMIC_CMPSWAP_RTN
     %6:vgpr_32 = FLAT_ATOMIC_CMPSWAP_RTN killed %7, killed %8, 0, 1, implicit $exec, implicit $flat_scr :: (load store seq_cst seq_cst (s32))
     %9:sreg_64_xexec = V_CMP_EQ_U32_e64 %6, %2, implicit $exec
     %10:vgpr_32 = V_CNDMASK_B32_e64 0, 0, 0, 1, killed %9, implicit $exec
     $vgpr0 = COPY %6
     $vgpr1 = COPY %10
     SI_RETURN implicit $vgpr0, implicit $vgpr1
 ...
 
 ---
 name:            atomic_inc
 tracksRegLiveness: true
 body:             |
   bb.0:
   ; CHECK-LABEL: MachineUniformityInfo for function: atomic_inc
     %2:vgpr_32 = IMPLICIT_DEF
     %1:vgpr_32 = IMPLICIT_DEF
     %0:vgpr_32 = IMPLICIT_DEF
     %3:sreg_64 = REG_SEQUENCE %0, %subreg.sub0, %1, %subreg.sub1
     %5:vreg_64 = COPY %3
     ; CHECK: DIVERGENT
     ; CHECK-SAME: GLOBAL_ATOMIC_INC_RTN
     %4:vgpr_32 = GLOBAL_ATOMIC_INC_RTN killed %5, %2, 0, 1, implicit $exec :: (load store (s32), addrspace 1)
     $vgpr0 = COPY %4
     SI_RETURN implicit $vgpr0
 ...
 
 ---
 name:            atomic_inc_64
 tracksRegLiveness: true
 body:             |
   bb.0:
   ; CHECK-LABEL: MachineUniformityInfo for function: atomic_inc_64
     %3:vgpr_32 = IMPLICIT_DEF
     %2:vgpr_32 = IMPLICIT_DEF
     %1:vgpr_32 = IMPLICIT_DEF
     %0:vgpr_32 = IMPLICIT_DEF
     %4:sreg_64 = REG_SEQUENCE %2, %subreg.sub0, %3, %subreg.sub1
     %5:sreg_64 = REG_SEQUENCE %0, %subreg.sub0, %1, %subreg.sub1
     %7:vreg_64 = COPY %5
     %8:vreg_64 = COPY %4
     ; CHECK: DIVERGENT
     ; CHECK-SAME: GLOBAL_ATOMIC_INC_X2_RTN
     %6:vreg_64 = GLOBAL_ATOMIC_INC_X2_RTN killed %7, killed %8, 0, 1, implicit $exec :: (load store (s64), addrspace 1)
     %9:vgpr_32 = COPY %6.sub1
     %10:vgpr_32 = COPY %6.sub0
     $vgpr0 = COPY %10
     $vgpr1 = COPY %9
     SI_RETURN implicit $vgpr0, implicit $vgpr1
 ...
 
 ---
 name:            atomic_dec
 tracksRegLiveness: true
 body:             |
   bb.0:
   ; CHECK-LABEL: MachineUniformityInfo for function: atomic_dec
     %2:vgpr_32 = IMPLICIT_DEF
     %1:vgpr_32 = IMPLICIT_DEF
     %0:vgpr_32 = IMPLICIT_DEF
     %3:sreg_64 = REG_SEQUENCE %0, %subreg.sub0, %1, %subreg.sub1
     %5:vreg_64 = COPY %3
     ; CHECK: DIVERGENT
     ; CHECK-SAME: GLOBAL_ATOMIC_DEC_RTN
     %4:vgpr_32 = GLOBAL_ATOMIC_DEC_RTN killed %5, %2, 0, 1, implicit $exec :: (load store (s32), addrspace 1)
     $vgpr0 = COPY %4
     SI_RETURN implicit $vgpr0
 ...
 
 
 ---
 name:            atomic_dec_64
 tracksRegLiveness: true
 body:             |
   bb.0:
   ; CHECK-LABEL: MachineUniformityInfo for function: atomic_dec_64
     %3:vgpr_32 = IMPLICIT_DEF
     %2:vgpr_32 = IMPLICIT_DEF
     %1:vgpr_32 = IMPLICIT_DEF
     %0:vgpr_32 = IMPLICIT_DEF
     %4:sreg_64 = REG_SEQUENCE %2, %subreg.sub0, %3, %subreg.sub1
     %5:sreg_64 = REG_SEQUENCE %0, %subreg.sub0, %1, %subreg.sub1
     %7:vreg_64 = COPY %5
     %8:vreg_64 = COPY %4
     ; CHECK: DIVERGENT
     ; CHECK-SAME: GLOBAL_ATOMIC_DEC_X2_RTN
     %6:vreg_64 = GLOBAL_ATOMIC_DEC_X2_RTN killed %7, killed %8, 0, 1, implicit $exec :: (load store (s64), addrspace 1)
     %9:vgpr_32 = COPY %6.sub1
     %10:vgpr_32 = COPY %6.sub0
     $vgpr0 = COPY %10
     $vgpr1 = COPY %9
     SI_RETURN implicit $vgpr0, implicit $vgpr1
 ...

llvm/test/Analysis/UniformityAnalysis/AMDGPU/MIR/hidden-diverge-gmir.mir

 # RUN: llc -mtriple=amdgcn-- -run-pass=print-machine-uniformity -o - %s 2>&1 | FileCheck %s
 # CHECK-LABEL: MachineUniformityInfo for function: hidden_diverge
 # CHECK-LABEL: BLOCK bb.0
 # CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.workitem.id.x)
 # CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s1) = G_ICMP intpred(slt)
 # CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s1) = G_XOR %{{[0-9]*}}:_, %{{[0-9]*}}:_
 # CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s1), %{{[0-9]*}}:_(s64) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.amdgcn.if)
 # CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s1), %{{[0-9]*}}:_(s64) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.amdgcn.if)
 # CHECK: DIVERGENT: G_BRCOND %{{[0-9]*}}:_(s1), %bb.1
 # CHECK: DIVERGENT: G_BR %bb.2
 # CHECK-LABEL: BLOCK bb.1
 # CHECK-LABEL: BLOCK bb.2
 # CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s32) = G_PHI %{{[0-9]*}}:_(s32), %bb.1, %{{[0-9]*}}:_(s32), %bb.0
 # CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s1) = G_PHI %{{[0-9]*}}:_(s1), %bb.1, %{{[0-9]*}}:_(s1), %bb.0
 # CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s1), %{{[0-9]*}}:_(s64) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.amdgcn.if)
 # CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s1), %{{[0-9]*}}:_(s64) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.amdgcn.if)
 # CHECK: DIVERGENT: G_BRCOND %{{[0-9]*}}:_(s1), %bb.3
 # CHECK: DIVERGENT: G_BR %bb.4
 # CHECK-LABEL: BLOCK bb.3
 # CHECK-LABEL: BLOCK bb.4
 # CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s32) = G_PHI %{{[0-9]*}}:_(s32), %bb.2, %{{[0-9]*}}:_(s32), %bb.3
 
 ---
 name:            hidden_diverge
 tracksRegLiveness: true
 body:             |
   bb.1:
     successors: %bb.2(0x40000000), %bb.3(0x40000000)
     liveins: $sgpr4_sgpr5
 
     %4:_(p4) = COPY $sgpr4_sgpr5
     %15:_(s32) = G_CONSTANT i32 0
     %17:_(s1) = G_CONSTANT i1 true
     %23:_(s32) = G_CONSTANT i32 1
     %30:_(s32) = G_CONSTANT i32 2
     %32:_(p1) = G_IMPLICIT_DEF
     %33:_(s32) = G_IMPLICIT_DEF
     %8:_(p4) = G_INTRINSIC intrinsic(@llvm.amdgcn.kernarg.segment.ptr)
     %9:_(<3 x s32>) = G_LOAD %8(p4) :: (dereferenceable invariant load (<3 x s32>), align 16, addrspace 4)
     %10:_(s64) = G_CONSTANT i64 4
     %11:_(p4) = G_PTR_ADD %8, %10(s64)
     %12:_(s64) = G_CONSTANT i64 8
     %13:_(p4) = G_PTR_ADD %8, %12(s64)
     %14:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.workitem.id.x)
     %16:_(s1) = G_ICMP intpred(slt), %14(s32), %15
     %18:_(s1) = G_XOR %16, %17
     %19:_(s1), %20:_(s64) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.amdgcn.if), %16(s1)
     G_BRCOND %19(s1), %bb.2
     G_BR %bb.3
 
   bb.2:
     successors: %bb.3(0x80000000)
 
     %21:_(s32) = G_EXTRACT_VECTOR_ELT %9(<3 x s32>), %15(s32)
     %22:_(s32) = G_EXTRACT_VECTOR_ELT %9(<3 x s32>), %23(s32)
     %24:_(s1) = G_ICMP intpred(slt), %21(s32), %15
 
   bb.3:
     successors: %bb.4(0x40000000), %bb.5(0x40000000)
 
     %25:_(s32) = G_PHI %22(s32), %bb.2, %33(s32), %bb.1
     %26:_(s1) = G_PHI %24(s1), %bb.2, %18(s1), %bb.1
     G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.amdgcn.end.cf), %20(s64)
     %27:_(s1), %28:_(s64) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.amdgcn.if), %26(s1)
     G_BRCOND %27(s1), %bb.4
     G_BR %bb.5
 
   bb.4:
     successors: %bb.5(0x80000000)
 
     %29:_(s32) = G_EXTRACT_VECTOR_ELT %9(<3 x s32>), %30(s32)
 
   bb.5:
     %31:_(s32) = G_PHI %25(s32), %bb.3, %29(s32), %bb.4
     G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.amdgcn.end.cf), %28(s64)
     G_STORE %31(s32), %32(p1) :: (volatile store (s32) into `i32 addrspace(1)* undef`, addrspace 1)
     S_ENDPGM 0
 
 ...

llvm/test/Analysis/UniformityAnalysis/AMDGPU/MIR/hidden-diverge.mir

 # RUN: llc -mtriple=amdgcn-- -run-pass=print-machine-uniformity -o - %s 2>&1 | FileCheck %s
 
 ---
 # CHECK-LABEL: MachineUniformityInfo for function: hidden_diverge
 # CHECK-LABEL: BLOCK bb.0
 # CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:vgpr_32(s32) = COPY $vgpr0
 # CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:sreg_64 = V_CMP_GT_I32_e64
 # CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:sreg_64 = V_CMP_LT_I32_e64
 # CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:vreg_1 = COPY
 # CHECK: DIVERGENT: %{{[0-9]*}}:sreg_64 = SI_IF
 # CHECK: DIVERGENT: S_BRANCH %bb.1
 # CHECK-LABEL: BLOCK bb.2
 # CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:sreg_32 = PHI %{{[0-9]*}}:sreg_32, %bb.0, %{{[0-9]*}}:sreg_32, %bb.1
 # CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:vreg_1 = PHI %{{[0-9]*}}:vreg_1, %bb.0, %{{[0-9]*}}:sreg_64, %bb.1
 # CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:sreg_64 = COPY %{{[0-9]*}}:vreg_1
 # CHECK: DIVERGENT: %{{[0-9]*}}:sreg_64 = SI_IF %{{[0-9]*}}:sreg_64, %bb.4
 # CHECK: DIVERGENT: S_BRANCH %bb.3
 # CHECK-LABEL: BLOCK bb.3
 # CHECK-LABEL: BLOCK bb.4
 # CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:vgpr_32 = PHI %{{[0-9]*}}:sreg_32, %bb.2, %{{[0-9]*}}:sreg_32, %bb.3
 
 name:            hidden_diverge
 tracksRegLiveness: true
 body:             |
   bb.0:
     successors: %bb.1(0x40000000), %bb.2(0x40000000)
     liveins: $vgpr0, $sgpr0_sgpr1
 
     %11:sgpr_64(p4) = COPY $sgpr0_sgpr1
     %10:vgpr_32(s32) = COPY $vgpr0
     %15:sreg_64_xexec = S_LOAD_DWORDX2_IMM %11(p4), 36, 0
     %16:sreg_32_xm0_xexec = S_LOAD_DWORD_IMM %11(p4), 44, 0
     %17:sreg_32 = COPY %15.sub1
     %18:sreg_32 = COPY %15.sub0
     %19:sgpr_96 = REG_SEQUENCE killed %18, %subreg.sub0, killed %17, %subreg.sub1, killed %16, %subreg.sub2
     %0:sgpr_96 = COPY %19
     %20:sreg_32 = S_MOV_B32 -1
     %21:sreg_64 = V_CMP_GT_I32_e64 %10(s32), killed %20, implicit $exec
     %22:sreg_32 = S_MOV_B32 0
     %23:sreg_64 = V_CMP_LT_I32_e64 %10(s32), killed %22, implicit $exec
     %1:vreg_1 = COPY %21
     %14:sreg_32 = IMPLICIT_DEF
     %2:sreg_64 = SI_IF killed %23, %bb.2, implicit-def dead $exec, implicit-def dead $scc, implicit $exec
     S_BRANCH %bb.1
 
   bb.1:
     successors: %bb.2(0x80000000)
 
     %24:sreg_32 = COPY %0.sub0
     %3:sreg_32 = COPY %0.sub1
     %25:sreg_32 = S_MOV_B32 0
     S_CMP_LT_I32 killed %24, killed %25, implicit-def $scc
     %26:sreg_64 = COPY $scc
     %4:sreg_64 = COPY %26
 
   bb.2:
     successors: %bb.3(0x40000000), %bb.4(0x40000000)
 
     %5:sreg_32 = PHI %14, %bb.0, %3, %bb.1
     %6:vreg_1 = PHI %1, %bb.0, %4, %bb.1
     SI_END_CF %2, implicit-def dead $exec, implicit-def dead $scc, implicit $exec
     %27:sreg_64 = COPY %6
     %7:sreg_64 = SI_IF %27, %bb.4, implicit-def dead $exec, implicit-def dead $scc, implicit $exec
     S_BRANCH %bb.3
 
   bb.3:
     successors: %bb.4(0x80000000)
 
     %8:sreg_32 = COPY %0.sub2
 
   bb.4:
     %9:vgpr_32 = PHI %5, %bb.2, %8, %bb.3
     SI_END_CF %7, implicit-def dead $exec, implicit-def dead $scc, implicit $exec
     %28:sreg_64 = IMPLICIT_DEF
     %29:vreg_64 = COPY %28
     GLOBAL_STORE_DWORD killed %29, %9, 0, 0, implicit $exec
     S_ENDPGM 0
 
 ...

llvm/test/Analysis/UniformityAnalysis/AMDGPU/MIR/irreducible/branch-outside-gmir.mir

 # RUN: llc -mtriple=amdgcn-- -run-pass=print-machine-uniformity -o - %s 2>&1 | FileCheck %s
 
 # CHECK-LABEL: MachineUniformityInfo for function: basic
 # CHECK-NEXT: CYCLES ASSSUMED DIVERGENT:
 # CHECK-NEXT: depth=1: entries(bb.1 bb.3) bb.2
 # CHECK-LABEL: BLOCK bb.1
 # CHECK: DIVERGENT
 # CHECK: DIVERGENT
 # CHECK-LABEL: BLOCK bb.2
 # CHECK: DIVERGENT
 # CHECK: DIVERGENT
 # CHECK-LABEL: BLOCK bb.3
 # CHECK: DIVERGENT
 # CHECK: DIVERGENT
 # CHECK-LABEL: BLOCK bb.4
 # CHECK-NOT: DIVERGENT
 
 
 ---
 name:            basic
 tracksRegLiveness: true
 body:             |
   bb.0:
     successors: %bb.3, %bb.1
 
     %0:_(s32) = G_IMPLICIT_DEF
     %1:_(s32) = G_CONSTANT i32 0
     %2:_(s32) = G_IMPLICIT_DEF
     %3:_(s32) = G_CONSTANT i32 1
     %4:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.workitem.id.x)
     %6:_(s1) = G_ICMP intpred(slt), %1(s32), %0(s32) ;uniform condition
     %7:_(s1) = G_ICMP intpred(eq), %4(s32), %0 ;divergent condition
     G_BRCOND %7(s1), %bb.3
     G_BR %bb.1
 
   bb.1:
     successors: %bb.2
 
     %8:_(s32) = G_PHI %0(s32), %bb.0, %2(s32), %bb.3
     %9:_(s32) = G_ADD %8(s32), %3(s32)
     G_BR %bb.2
 
   bb.2:
     successors: %bb.3, %bb.4
 
     %13:_(s32) = G_ADD %2(s32), %3(s32)
     %10:_(s32) = G_ADD %8(s32), %3(s32)
     G_BRCOND %6(s1), %bb.3
     G_BR %bb.4
 
   bb.3:
     successors: %bb.1
     %11:_(s32) = G_PHI %13(s32), %bb.2, %0(s32), %bb.0
     %12:_(s32) = G_ADD %11(s32), %3(s32)
     G_BR %bb.1
   bb.4:
     %14:_(s32) = G_ADD %2(s32), %3(s32)
     S_ENDPGM 0
 ...

llvm/test/Analysis/UniformityAnalysis/AMDGPU/MIR/irreducible/diverged-entry-basic-gmir.mir

 # RUN: llc -mtriple=amdgcn-- -run-pass=print-machine-uniformity -o - %s 2>&1 | FileCheck %s
 # CHECK-LABEL: MachineUniformityInfo for function: divergent_cycle_1
 # CHECK-NEXT: CYCLES ASSSUMED DIVERGENT:
 # CHECK-NEXT: depth=1: entries(bb.3 bb.1) bb.4 bb.2
 # CHECK-NEXT: CYCLES WITH DIVERGENT EXIT:
 # CHECK-NEXT: depth=2: entries(bb.4 bb.1) bb.2
 # CHECK-NEXT: depth=1: entries(bb.3 bb.1) bb.4 bb.2
 
 
 
 ---
 name:            divergent_cycle_1
 tracksRegLiveness: true
 body:             |
   bb.0:
     successors: %bb.1, %bb.3
     %0:_(s32) = G_CONSTANT i32 0
     %1:_(s32) = G_CONSTANT i32 1
 
     %2:_(s32) = G_IMPLICIT_DEF
     %3:_(s32) = G_IMPLICIT_DEF
 
     %4:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.workitem.id.x)
     %6:_(s1) = G_ICMP intpred(slt), %2(s32), %0(s32) ;uniform condition
     %7:_(s1) = G_ICMP intpred(eq), %4(s32), %0 ;divergent condition
     G_BRCOND %6(s1), %bb.1
     G_BR %bb.3
 
   bb.1:
     successors: %bb.2
 
   ; CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s32) = G_PHI %{{[0-9]*}}:_(s32), %bb.0, %{{[0-9]*}}:_(s32), %bb.4
     %8:_(s32) = G_PHI %2(s32), %bb.0, %3(s32), %bb.4
     %9:_(s32) = G_ADD %3(s32), %1(s32)
     G_BR %bb.2
 
   bb.2:
     successors: %bb.3, %bb.4
 
     %13:_(s32) = G_ADD %3(s32), %1(s32)
     G_BRCOND %7(s1), %bb.4
     G_BR %bb.3
 
   bb.3:
     successors: %bb.4
 
     %14:_(s32) = G_ADD %3(s32), %1(s32)
     G_BR %bb.4
   bb.4:
     successors: %bb.5, %bb.1
 
   ; CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s32) = G_PHI %{{[0-9]*}}:_(s32), %bb.2, %{{[0-9]*}}:_(s32), %bb.3
     %15:_(s32) = G_PHI %13(s32), %bb.2, %14(s32), %bb.3
     %16:_(s32) = G_ADD %3(s32), %1(s32)
     G_BRCOND %6(s1), %bb.5
     G_BR %bb.1
 
   bb.5:
     %17:_(s32) = G_ADD %3(s32), %1(s32)
     S_ENDPGM 0
 ...
 
 # CHECK-LABEL: MachineUniformityInfo for function: uniform_cycle_1
 ---
 name:            uniform_cycle_1
 tracksRegLiveness: true
 body:             |
   bb.0:
     successors: %bb.1, %bb.5
     %0:_(s32) = G_CONSTANT i32 0
     %1:_(s32) = G_CONSTANT i32 1
 
     %2:_(s32) = G_IMPLICIT_DEF
     %3:_(s32) = G_IMPLICIT_DEF
 
     %4:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.workitem.id.x)
     %6:_(s1) = G_ICMP intpred(slt), %2(s32), %0(s32) ;uniform condition
     %7:_(s1) = G_ICMP intpred(eq), %4(s32), %0 ;divergent condition
     G_BRCOND %6(s1), %bb.1
     G_BR %bb.5
 
   bb.1:
     successors: %bb.2
 
   ; CHECK-NOT: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s32) = G_PHI %{{[0-9]*}}:_(s32), %bb.0, %{{[0-9]*}}:_(s32), %bb.4
     %8:_(s32) = G_PHI %2(s32), %bb.0, %3(s32), %bb.5
     %9:_(s32) = G_ADD %3(s32), %1(s32)
     G_BR %bb.2
 
   bb.2:
     successors: %bb.3, %bb.4
 
     %13:_(s32) = G_ADD %3(s32), %1(s32)
     G_BRCOND %7(s1), %bb.4
     G_BR %bb.3
 
   bb.3:
     successors: %bb.4
 
     %14:_(s32) = G_ADD %3(s32), %1(s32)
     G_BR %bb.4
   bb.4:
     successors: %bb.6, %bb.5
 
   ; CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s32) = G_PHI %{{[0-9]*}}:_(s32), %bb.2, %{{[0-9]*}}:_(s32), %bb.3
     %15:_(s32) = G_PHI %13(s32), %bb.2, %14(s32), %bb.3
     %16:_(s32) = G_ADD %3(s32), %1(s32)
     G_BRCOND %6(s1), %bb.6
     G_BR %bb.5
 
   bb.5:
     successors: %bb.1
     %18:_(s32) = G_ADD %3(s32), %1(s32)
     G_BR %bb.1
   bb.6:
     %17:_(s32) = G_ADD %3(s32), %1(s32)
     S_ENDPGM 0
 ...

llvm/test/Analysis/UniformityAnalysis/AMDGPU/MIR/irreducible/exit-divergence-gmir.mir

 # RUN: llc -mtriple=amdgcn-- -run-pass=print-machine-uniformity -o - %s 2>&1 | FileCheck %s
 # CHECK-LABEL: MachineUniformityInfo for function: basic
 # CHECK-NOT: CYCLES ASSSUMED DIVERGENT:
 # CHECK: CYCLES WITH DIVERGENT EXIT:
 # CHECK: depth=1: entries(bb.1 bb.3) bb.2
 
 ---
 name:            basic
 tracksRegLiveness: true
 body:             |
   bb.0:
     successors: %bb.3, %bb.1
 
     %0:_(s32) = G_CONSTANT i32 0
     %1:_(s32) = G_CONSTANT i32 1
 
     %2:_(s32) = G_IMPLICIT_DEF
     %3:_(s32) = G_IMPLICIT_DEF
 
     %6:_(s1) = G_ICMP intpred(slt), %2(s32), %0(s32) ;uniform condition
     %7:_(s1) = G_ICMP intpred(eq), %4(s32), %0 ;divergent condition
     %4:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.workitem.id.x)
     G_BRCOND %6(s1), %bb.3
     G_BR %bb.1
 
   bb.1:
     successors: %bb.2
 
     %8:_(s32) = G_PHI %2(s32), %bb.0, %3(s32), %bb.3
     %10:_(s32) = G_PHI %2(s32), %bb.0, %16(s32), %bb.3
     %9:_(s32) = G_ADD %3(s32), %1(s32)
     G_BR %bb.2
 
   bb.2:
     successors: %bb.3, %bb.4
 
     %13:_(s32) = G_ADD %3(s32), %1(s32)
     %14:_(s32) = G_ADD %10(s32), %1(s32)
     %15:_(s32) = G_ADD %10(s32), %1(s32)
 
     G_BRCOND %7(s1), %bb.3
     G_BR %bb.4
 
   bb.3:
     successors: %bb.1
     %16:_(s32) = G_PHI %13(s32), %bb.2, %2(s32), %bb.0
     %17:_(s32) = G_ADD %3(s32), %1(s32)
     G_BR %bb.1
   bb.4:
     ; CHECK-LABEL: bb.4
     ; CHECK: DIVERGENT:
     ; CHECK: DIVERGENT:
     ; CHECK-NOT: DIVERGENT:
     %18:_(s32) = G_ADD %8(s32), %3(s32)
     %19:_(s32) = G_ADD %8(s32), %3(s32)
     %20:_(s32) = G_ADD %3(s32), %1(s32)
     S_ENDPGM 0
 ...

llvm/test/Analysis/UniformityAnalysis/AMDGPU/MIR/irreducible/irreducible-1.mir

 # RUN: llc -mtriple=amdgcn-- -run-pass=print-machine-uniformity -o - %s 2>&1 | FileCheck %s
 
 # CHECK-LABEL: MachineUniformityInfo for function: irreducible
 # CHECK: CYCLES ASSSUMED DIVERGENT:
 # CHECK:   depth=1: entries(bb.2 bb.1) bb.3 bb.5 bb.4
 # CHECK: CYCLES WITH DIVERGENT EXIT:
 # CHECK-DAG:   depth=1: entries(bb.2 bb.1) bb.3 bb.5 bb.4
 # CHECK-DAG:   depth=2: entries(bb.3 bb.1) bb.5 bb.4
 
 ---
 name:            irreducible
 tracksRegLiveness: true
 machineFunctionInfo:
   isEntryFunction: true
 body:             |
   bb.0:
     successors: %bb.1, %bb.2
     liveins: $vgpr0, $vgpr1, $vgpr2, $sgpr4_sgpr5, $sgpr6_sgpr7, $sgpr8_sgpr9, $sgpr10_sgpr11, $sgpr14, $sgpr15, $sgpr16
 
     %0:sreg_32 = IMPLICIT_DEF
     %2:vgpr_32 = COPY $vgpr0
     %3:vgpr_32 = V_MOV_B32_e32 0, implicit $exec
     S_CMP_EQ_U32 %0, 0, implicit-def $scc
     S_CBRANCH_SCC1 %bb.1, implicit $scc
     S_BRANCH %bb.2
 
   bb.1:
     %28:vgpr_32 = PHI %3, %bb.0, %49, %bb.5
     %29:vgpr_32 = V_ADD_U32_e64 %28, 1, 0, implicit $exec
     S_BRANCH %bb.3
 
   bb.2:
     %38:vgpr_32 = PHI %3, %bb.0, %49, %bb.4
     %39:vgpr_32 = V_ADD_U32_e64 %38, 2, 0, implicit $exec
 
   bb.3:
     %49:vgpr_32 = PHI %29, %bb.1, %39, %bb.2
 
   bb.4:
     successors: %bb.2, %bb.5
 
     %50:vgpr_32 = V_AND_B32_e32 3, %2, implicit $exec
     %51:sreg_64 = V_CMP_EQ_U32_e64 %50, 2, implicit $exec
     %52:sreg_64 = SI_IF killed %51:sreg_64, %bb.2, implicit-def dead $exec, implicit-def dead $scc, implicit $exec
 
   bb.5:
     successors: %bb.1, %bb.6
     %61:sreg_64 = V_CMP_EQ_U32_e64 %50, 1, implicit $exec
     %62:sreg_64 = SI_IF killed %61:sreg_64, %bb.1, implicit-def dead $exec, implicit-def dead $scc, implicit $exec
 
   bb.6:
     S_ENDPGM 0
 ...

llvm/test/Analysis/UniformityAnalysis/AMDGPU/MIR/irreducible/irreducible-2-gmir.mir

 # RUN: llc -mtriple=amdgcn-- -run-pass=print-machine-uniformity -o - %s 2>&1 | FileCheck %s
 
 #        bb0(div)
 #       /    \
 #     bb1 <-> bb2
 #              |
 #             bb3
 # CHECK-LABEL: MachineUniformityInfo for function: cycle_diverge_enter
 # CHECK-NEXT: CYCLES ASSSUMED DIVERGENT:
 # CHECK-NEXT: depth=1: entries(bb.2 bb.1)
 # CHECK-NEXT: CYCLES WITH DIVERGENT EXIT:
 # CHECK-NEXT: depth=1: entries(bb.2 bb.1)
 ---
 name:            cycle_diverge_enter
 tracksRegLiveness: true
 body:             |
   bb.0:
     successors: %bb.1, %bb.2
 
     %0:_(s32) = G_IMPLICIT_DEF
     %1:_(s32) = G_CONSTANT i32 0
     %2:_(s32) = G_IMPLICIT_DEF
     %3:_(s32) = G_CONSTANT i32 1
     %4:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.workitem.id.x)
     %6:_(s1) = G_ICMP intpred(slt), %4(s32), %1 ; DIVERGENT CONDITION
     %7:_(s1) = G_ICMP intpred(slt), %2(s32), %1 ; UNIFORM CONDITION
     G_BRCOND %6(s1), %bb.1 ; divergent branch
     G_BR %bb.2
   bb.1:
     successors: %bb.2
 
     %8:_(s32) = G_PHI %1(s32), %bb.0, %0(s32), %bb.2
     G_BR %bb.2
   bb.2:
     successors: %bb.1, %bb.3
 
     %9:_(s32) = G_PHI %2(s32), %bb.1, %3(s32), %bb.0
     %10:_(s1) = G_ICMP intpred(eq), %9(s32), %1(s32)
     G_BRCOND %10(s1), %bb.3  ; divergent branch
     G_BR %bb.1
 
   bb.3:
     %11:_(s32), %12:_(s1) = G_UADDO %9, %3
     S_ENDPGM 0
 
 ...
 
 
 # CHECK-LABEL: MachineUniformityInfo for function: cycle_diverge_exit
 # CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s32), %{{[0-9]*}}:_(s1) = G_UADDO %8:_, %{{[0-9]*}}:_
 #        bb0
 #       /    \
 #     bb1 <-> bb2(div)
 #              |
 #             bb3
 ---
 name:            cycle_diverge_exit
 tracksRegLiveness: true
 body:             |
   bb.0:
     successors: %bb.1, %bb.2
 
     %0:_(s32) = G_IMPLICIT_DEF
     %1:_(s32) = G_CONSTANT i32 0
     %2:_(s32) = G_IMPLICIT_DEF
     %3:_(s32) = G_CONSTANT i32 1
     %4:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.workitem.id.x)
     %6:_(s1) = G_ICMP intpred(slt), %4(s32), %1 ; DIVERGENT CONDITION
     %7:_(s1) = G_ICMP intpred(slt), %2(s32), %1 ; UNIFORM CONDITION
     G_BRCOND %7(s1), %bb.1 ; uniform branch
     G_BR %bb.2
   bb.1:
     successors: %bb.2
 
     %8:_(s32) = G_PHI %1(s32), %bb.0, %0(s32), %bb.2
     G_BR %bb.2
   bb.2:
     successors: %bb.1, %bb.3
 
     %9:_(s32) = G_PHI %2(s32), %bb.1, %3(s32), %bb.0
     %10:_(s1) = G_ICMP intpred(sgt), %9(s32), %1(s32)
     G_BRCOND %6(s1), %bb.3  ; divergent branch
     G_BR %bb.1
 
   bb.3:
     %11:_(s32), %12:_(s1) = G_UADDO %9, %3
     S_ENDPGM 0
 
 ...

llvm/test/Analysis/UniformityAnalysis/AMDGPU/MIR/join-loopexit-gmir.mir

 # RUN: llc -mtriple=amdgcn-- -run-pass=print-machine-uniformity -o - %s 2>&1 | FileCheck %s
 # CHECK-LABEL: MachineUniformityInfo for function: test
 
 # CHECK-LABEL: BLOCK bb.0
 # CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s1) = G_ICMP intpred(eq), %{{[0-9]*}}:_(s32), %{{[0-9]*}}:_
 
 # CHECK-LABEL: BLOCK bb.1
 # CHECK-NOT: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s32) = G_PHI %{{[0-9]*}}:_(s32), %bb.0, %{{[0-9]*}}:_(s32), %bb.2
 # CHECK:  DIVERGENT: G_BRCOND %{{[0-9]*}}:_(s1), %bb.3
 
 # CHECK-LABEL: BLOCK bb.2
 # CHECK-NOT: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s32), %{{[0-9]*}}:_(s1) = G_UADDO_
 # CHECK-NOT: DIVERGENT: G_BRCOND %{{[0-9]*}}:_(s1), %bb.3
 
 # CHECK-LABEL: BLOCK bb.3
 # CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s1) = G_PHI %{{[0-9]*}}:_(s1), %bb.1, %{{[0-9]*}}:_(s1), %bb.2
 # CHECK: DIVERGENT: G_BRCOND %{{[0-9]*}}:_(s1), %bb.4
 # CHECK: DIVERGENT: G_BR %bb.5
 
 ---
 name:            test
 tracksRegLiveness: true
 body:             |
   bb.0:
     successors: %bb.1
 
     %2:_(s1) = G_CONSTANT i1 true
     %3:_(s1) = G_CONSTANT i1 false
     %1:_(s32) = G_CONSTANT i32 0
     %20:_(s32) = G_CONSTANT i32 7
     %5:_(s32) = G_CONSTANT i32 -1
     %4:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.workitem.id.x)
     %6:_(s1) = G_ICMP intpred(eq), %4(s32), %5
 
   bb.1:
     successors: %bb.2, %bb.3
 
     %8:_(s32) = G_PHI %20(s32), %bb.0, %21(s32), %bb.2
     G_BRCOND %6(s1), %bb.3   ; Entrance to loop is divergent
   bb.2:
     successors: %bb.3, %bb.1
 
     %21:_(s32), %22:_(s1) = G_UADDO %8, %5
     %23:_(s1) = G_ICMP intpred(eq), %21(s32), %1
     G_BRCOND %23(s1), %bb.3
     G_BR %bb.1
 
   bb.3:
     %31:_(s1) = G_PHI %2(s1), %bb.1, %3(s1), %bb.2
     S_ENDPGM 0
     G_BRCOND %31(s1), %bb.4
     G_BR %bb.5
   bb.4:
     G_BR %bb.5
   bb.5:
 
 ...

llvm/test/Analysis/UniformityAnalysis/AMDGPU/MIR/loads-gmir.mir

 # RUN: llc -mtriple=amdgcn-- -run-pass=print-machine-uniformity -o - %s 2>&1 | FileCheck %s
 
 ---
 name:            loads
 tracksRegLiveness: true
 body:             |
   bb.1.entry:
     %1:_(p0) = G_IMPLICIT_DEF
     %4:_(p1) = G_IMPLICIT_DEF
     %6:_(p5) = G_IMPLICIT_DEF
 
     ; Atomic load
     ; CHECK: DIVERGENT
     ; CHECK-SAME: G_LOAD
     %0:_(s32) = G_LOAD %1(p0) :: (load seq_cst (s32) from `ptr undef`)
 
     ; flat load
     ; CHECK: DIVERGENT
     ; CHECK-SAME: G_LOAD
     %2:_(s32) = G_LOAD %1(p0) :: (load (s32) from `ptr undef`)
 
     ; Gloabal load
     ; CHECK-NOT: DIVERGENT
     %3:_(s32) = G_LOAD %4(p1) :: (load (s32) from `ptr addrspace(1) undef`, addrspace 1)
 
     ; Private load
     ; CHECK: DIVERGENT
     ; CHECK-SAME: G_LOAD
     %5:_(s32) = G_LOAD %6(p5) :: (volatile load (s32) from `ptr addrspace(5) undef`, addrspace 5)
     G_STORE %2(s32), %4(p1) :: (volatile store (s32) into `ptr addrspace(1) undef`, addrspace 1)
     G_STORE %3(s32), %4(p1) :: (volatile store (s32) into `ptr addrspace(1) undef`, addrspace 1)
     G_STORE %5(s32), %4(p1) :: (volatile store (s32) into `ptr addrspace(1) undef`, addrspace 1)
     G_STORE %0(s32), %4(p1) :: (volatile store (s32) into `ptr addrspace(1) undef`, addrspace 1)
     SI_RETURN
 
 ...

llvm/test/Analysis/UniformityAnalysis/AMDGPU/MIR/never-uniform.mir

 # RUN: llc -mtriple=amdgcn-- -mcpu=gfx900 -run-pass=print-machine-uniformity -o - %s 2>&1 | FileCheck %s
 # loads from flat non uniform
 ---
 name:            flatloads
 tracksRegLiveness: true
 machineFunctionInfo:
   isEntryFunction: true
 
 body:             |
   bb.0:
     ; CHECK-LABEL: MachineUniformityInfo for function: flatloads
     ; CHECK: DIVERGENT: %1
     ; CHECK-NOT: DIVERGENT: %2
     %0:vreg_64 = IMPLICIT_DEF
     %1:vgpr_32(s32) = FLAT_LOAD_DWORD %0, 0, 0, implicit $exec, implicit $flat_scr :: (load (s32))
     %2:vgpr_32(s32) = FLAT_LOAD_DWORD %0, 0, 0, implicit $exec, implicit $flat_scr :: (load (s32), addrspace 1)
     %3:sreg_32 = V_READFIRSTLANE_B32 %1(s32), implicit $exec
     S_ENDPGM 0
 ...
 
 # loads from scratch non uniform
 ---
 name:            scratchloads
 tracksRegLiveness: true
 machineFunctionInfo:
   isEntryFunction: true
 
 body:             |
   bb.0:
     ; CHECK-LABEL: MachineUniformityInfo for function: scratchloads
     ; CHECK: DIVERGENT: %1
     %0:vgpr_32 = V_MOV_B32_e32 0, implicit $exec
     %1:vgpr_32 = SCRATCH_LOAD_DWORD %0, 0, 0, implicit $exec, implicit $flat_scr :: (load (s32), addrspace 5)
     S_ENDPGM 0
 ...
 
 # Global load
 ---
 name:            globalloads
 tracksRegLiveness: true
 machineFunctionInfo:
   isEntryFunction: true
 
 body:             |
   bb.0:
     ; CHECK-LABEL: MachineUniformityInfo for function: globalloads
     ; CHECK: DIVERGENT: %2
     ; CHECK-NOT: DIVERGENT: %3
     %0:vreg_64 = IMPLICIT_DEF
     %1:vreg_64 = IMPLICIT_DEF
     %2:vgpr_32(s32) = GLOBAL_LOAD_DWORD %0, 0, 0, implicit $exec, implicit $flat_scr :: (load (s32))
     %3:vreg_64 = GLOBAL_LOAD_DWORDX2 %1, 0, 0, implicit $exec :: (load (s64), addrspace 1)
     %4:sreg_32 = V_READFIRSTLANE_B32 %2(s32), implicit $exec
     S_ENDPGM 0
 ...
 
 # FIXME:: ADDTID might instruction incorrectly marked uniform
 ---
 name:            dsreads
 tracksRegLiveness: true
 machineFunctionInfo:
   isEntryFunction: true
 
 body:             |
   bb.0:
     ; CHECK-LABEL: MachineUniformityInfo for function: dsreads
     ; CHECK-NEXT: ALL VALUES UNIFORM
     %0:vreg_64 = IMPLICIT_DEF
     $m0 = S_MOV_B32 0
     %1:vgpr_32 = DS_READ_ADDTID_B32 0, 0, implicit $m0, implicit $exec
     S_ENDPGM 0
 ...
 
 # copy source == $sgpr => uniform, $vgpr => divergent
 ---
 name:            sgprcopy
 tracksRegLiveness: true
 machineFunctionInfo:
   isEntryFunction: true
 body:             |
   bb.0:
     ; CHECK-LABEL: MachineUniformityInfo for function: sgprcopy
     ; CHECK: DIVERGENT: %2
     liveins: $sgpr0,$sgpr1,$vgpr0
     %0:sgpr_32 = COPY $sgpr0
     %1:vgpr_32 = COPY $sgpr1
     %2:vgpr_32 = COPY $vgpr0
     S_ENDPGM 0
 ...
 
 # writelane is not uniform
 ---
 name:            writelane
 machineFunctionInfo:
   isEntryFunction: true
 body:             |
   bb.0:
     ; CHECK-LABEL: MachineUniformityInfo for function: writelane
     ; CHECK: DIVERGENT: %4
     ; CHECK: DIVERGENT: %5
     %0:vgpr_32 = IMPLICIT_DEF
     %1:vgpr_32 = IMPLICIT_DEF
     %2:sgpr_32 = V_READFIRSTLANE_B32 %0, implicit $exec
     %3:sgpr_32 = V_READLANE_B32 %1, 0, implicit $exec
     $sgpr0 = V_READFIRSTLANE_B32 $vgpr0, implicit $exec
     $sgpr1 = V_READLANE_B32 $vgpr1, $sgpr0, implicit $exec
 
     %4:vgpr_32 = V_WRITELANE_B32 0, 0, %0, implicit $exec
     %5:sreg_64 = V_CMP_EQ_U32_e64 %0, %4, implicit $exec
     S_CBRANCH_VCCZ %bb.1, implicit $vcc
 
   bb.1:
     %16:vgpr_32 = IMPLICIT_DEF
     S_ENDPGM 0
 ...
 # Direclty reading physing vgpr not uniform
 ---
 name:            physicalreg
 tracksRegLiveness: true
 body:             |
   bb.0:
     ; CHECK-LABEL: MachineUniformityInfo for function: physicalreg
     ; CHECK: DIVERGENT: %0
     ; CHECK: DIVERGENT: %1
     ; CHECK: DIVERGENT: %2
     ; CHECK: DIVERGENT: %3
     ; CHECK: DIVERGENT: %4
     ; CHECK-NOT: DIVERGENT
     ; CHECK: DIVERGENT: %5
     liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
     %0:vgpr_32 = COPY $vgpr0
     %1:vgpr_32 = COPY $vgpr1
     %2:vgpr_32 = V_AND_B32_e32 %1, $vgpr3, implicit $exec
     %3:vgpr_32 = V_ADD_U32_e32 $vgpr2, $vgpr3, implicit $exec
     %4:vgpr_32 = V_SUB_CO_U32_e32 $vgpr2, $vgpr3, implicit $exec, implicit-def $vcc
     %5:vgpr_32 = V_AND_B32_e32 $vgpr4, $vgpr5, implicit $exec
     S_ENDPGM 0
 ...

llvm/test/Analysis/UniformityAnalysis/AMDGPU/MIR/temporal-diverge-gmir.mir

 # RUN: llc -mtriple=amdgcn-- -run-pass=print-machine-uniformity -o - %s 2>&1 | FileCheck %s
 # CHECK-LABEL: MachineUniformityInfo for function: hidden_loop_diverge
 
 # CHECK-LABEL: BLOCK bb.0
 # CHECK-NOT: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s1) = G_ICMP intpred(slt), %{{[0-9]*}}:_(s32), %{{[0-9]*}}:_
 # CHECK-NOT: DIVERGENT: G_BRCOND %{{[0-9]*}}:_(s1), %bb.3
 # CHECK-NOT: DIVERGENT: G_BR %bb.1
 
 # CHECK-LABEL: BLOCK bb.1
 # CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s1) = G_ICMP intpred(slt), %{{[0-9]*}}:_(s32), %{{[0-9]*}}:_
 # CHECK: DIVERGENT: G_BRCOND %{{[0-9]*}}:_(s1), %bb.3
 # CHECK: DIVERGENT: G_BR %bb.2
 
 # CHECK-LABEL: BLOCK bb.2
 # CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s1) = G_ICMP intpred(sgt), %{{[0-9]*}}:_(s32), %{{[0-9]*}}:_
 # CHECK: DIVERGENT: G_BRCOND %{{[0-9]*}}:_(s1), %bb.4
 # CHECK: DIVERGENT: G_BR %bb.1
 
 # CHECK-LABEL: BLOCK bb.3
 # CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s32) = G_PHI %{{[0-9]*}}:_(s32), %bb.0, %{{[0-9]*}}:_(s32), %bb.1
 # CHECK-NOT: DIVERGENT: G_BRCOND %{{[0-9]*}}:_(s1), %bb.4
 # CHECK-NOT: DIVERGENT: G_BR %bb.5
 
 # CHECK-LABEL: BLOCK bb.4
 # CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s32) = G_PHI %{{[0-9]*}}:_(s32), %bb.3, %{{[0-9]*}}:_(s32), %bb.2
 
 # CHECK-LABEL: BLOCK bb.5
 # CHECK: DIVERGENT: %{{[0-9]*}}: %{{[0-9]*}}:_(s32) = G_PHI %{{[0-9]*}}:_(s32), %bb.3, %{{[0-9]*}}:_(s32), %bb.4
 
 ---
 name:            hidden_loop_diverge
 tracksRegLiveness: true
 body:             |
     bb.0:
         successors: %bb.3, %bb.1
         liveins: $sgpr4_sgpr5
 
         %0:_(s32) = G_IMPLICIT_DEF
         %20:_(s32) = G_IMPLICIT_DEF
         %21:_(s32) = G_CONSTANT i32 42
         %22:_(s32) = G_IMPLICIT_DEF
         %1:_(s32) = G_CONSTANT i32 0
         %2:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.workitem.id.x)
         %3:_(s1) = G_ICMP intpred(slt), %0(s32), %1
         G_BRCOND %3(s1), %bb.3                            ; Uniform branch
         G_BR %bb.1
     bb.1:
         successors: %bb.3, %bb.2
 
         %4:_(s32) = G_PHI %1(s32), %bb.0, %7(s32), %bb.2
         %5:_(s1) = G_ICMP intpred(slt), %1(s32), %2(s32)
         G_BRCOND %5(s1), %bb.3
         G_BR %bb.2
     bb.2:
         successors: %bb.4, %bb.1
 
         %6:_(s32) = G_CONSTANT i32 1
         %7:_(s32) = G_ADD %6(s32), %4(s32)
         %8:_(s1) = G_ICMP intpred(sgt), %2(s32), %1(s32)
         G_BRCOND %8(s1), %bb.4
         G_BR %bb.1
     bb.3:
         successors: %bb.4, %bb.5
 
         %9:_(s32) = G_PHI %20(s32), %bb.0, %4(s32), %bb.1    ; Temporal divergent phi
         G_BRCOND %3(s1), %bb.4
         G_BR %bb.5
 
     bb.4:
         successors: %bb.5
 
         %10:_(s32) = G_PHI %21(s32), %bb.3, %22(s32), %bb.2
         G_BR %bb.5
     bb.5:
         %11:_(s32) = G_PHI %20(s32), %bb.3, %22(s32), %bb.4
 ...

llvm/test/Analysis/UniformityAnalysis/AMDGPU/MIR/uses-value-from-cycle.mir

 # RUN: llc -mtriple=amdgcn-- -mcpu=gfx1030 -run-pass=print-machine-uniformity -o - %s 2>&1 | FileCheck %s
 ---
 name:            f1
 tracksRegLiveness: true
 body:             |
   ; CHECK-LABEL: MachineUniformityInfo for function: f1
   bb.1:
     %3:_(s32) = G_CONSTANT i32 0
     %25:_(s32) = G_IMPLICIT_DEF
   
   bb.2:
     %0:_(s32) = G_PHI %22(s32), %bb.5, %3(s32), %bb.1
     %1:_(s32) = G_PHI %3(s32), %bb.1, %20(s32), %bb.5
     %2:_(s32) = G_PHI %3(s32), %bb.1, %19(s32), %bb.5
     %36:_(s32) = G_CONSTANT i32 0
     %4:_(s1) = G_ICMP intpred(slt), %1(s32), %36
   
   bb.3:
     successors: %bb.4(0x04000000), %bb.3(0x7c000000)
   
     %5:_(s32) = G_PHI %38(s32), %bb.3, %1(s32), %bb.2
     %38:_(s32) = G_CONSTANT i32 0
     G_BRCOND %4(s1), %bb.3
     G_BR %bb.4
   
   bb.4:
     successors: %bb.7, %bb.5
   
     %6:_(s32) = G_PHI %5(s32), %bb.3
     %33:_(s1) = G_CONSTANT i1 true
     %7:_(s64) = G_SEXT %2(s32)
     %39:_(s32) = G_CONSTANT i32 2
     %10:_(s64) = G_SHL %7, %39(s32)
     %11:_(p0) = G_INTTOPTR %10(s64)
     %13:_(s32) = G_LOAD %11(p0) :: (load (s32))
     %37:_(s32) = G_CONSTANT i32 0
     %14:sreg_32_xm0_xexec(s1) = G_ICMP intpred(slt), %13(s32), %37
     %16:sreg_32_xm0_xexec(s32) = SI_IF %14(s1), %bb.5, implicit-def $exec, implicit-def $scc, implicit $exec
     G_BR %bb.7
   
   bb.5:
     successors: %bb.6(0x04000000), %bb.2(0x7c000000)
     ; CHECK: DIVERGENT: %{{[0-9]}}: %{{[0-9]}}:_(s32) = G_PHI
     ; CHECK: DIVERGENT: %{{[0-9]}}: %{{[0-9]}}:_(s32) = G_PHI
     ; CHECK: DIVERGENT: %{{[0-9]}}: %{{[0-9]}}:_(s32) = G_PHI
     ; CHECK-NOT: DIVERGENT: %{{[0-9]}}: %{{[0-9]}}:sreg_32_xm0_xexec(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.if.break)
     %19:_(s32) = G_PHI %18(s32), %bb.7, %25(s32), %bb.4
     %20:_(s32) = G_PHI %6(s32), %bb.7, %25(s32), %bb.4
     %21:_(s1) = G_PHI %34(s1), %bb.7, %33(s1), %bb.4
     G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.amdgcn.end.cf), %16(s32)
     %22:sreg_32_xm0_xexec(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.if.break), %21(s1), %0(s32)
     SI_LOOP %22(s32), %bb.2, implicit-def $exec, implicit-def $scc, implicit $exec
     G_BR %bb.6
   
   bb.6:
     %24:_(s32) = G_PHI %22(s32), %bb.5
     G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.amdgcn.end.cf), %24(s32)
     SI_RETURN
   
   bb.7:
     %34:_(s1) = G_CONSTANT i1 false
     %35:_(s32) = G_CONSTANT i32 1
     %18:_(s32) = G_OR %2, %35
     G_BR %bb.5
 
 ...

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

-; RUN: opt -mtriple amdgcn-unknown-amdhsa -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 ; RUN: opt -mtriple amdgcn-unknown-amdhsa -passes='print<uniformity>' -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)

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

-; RUN: opt -mtriple amdgcn-- -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 ; RUN: opt -mtriple amdgcn-- -passes='print<uniformity>' -disable-output %s 2>&1 | FileCheck %s
 
 ; CHECK: DIVERGENT: %orig = atomicrmw xchg ptr %ptr, i32 %val seq_cst
 define amdgpu_kernel void @test1(ptr %ptr, i32 %val) #0 {
   %orig = atomicrmw xchg ptr %ptr, i32 %val seq_cst
   store i32 %orig, ptr %ptr
   ret void
 }

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

-; RUN: opt -mtriple amdgcn-unknown-amdhsa -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 ; RUN: opt -mtriple amdgcn-unknown-amdhsa -passes='print<uniformity>' -disable-output %s 2>&1 | FileCheck %s
 
 declare i32 @gf2(i32)
 declare i32 @gf1(i32)
 
 define  void @tw1(ptr addrspace(4) noalias nocapture readonly %A, ptr addrspace(4) noalias nocapture %B) local_unnamed_addr #2 {
 ; CHECK: for function 'tw1':
 ; CHECK-DAG: DIVERGENT: ptr addrspace(4) %A

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

-; RUN: opt -mtriple amdgcn-mesa-mesa3d -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 ; RUN: opt -mtriple amdgcn-mesa-mesa3d -passes='print<uniformity>' -disable-output %s 2>&1 | FileCheck %s
 
 ; Tests control flow intrinsics that should be treated as uniform
 
 ; CHECK: for function 'test_if_break':
 ; CHECK: DIVERGENT: %cond = icmp eq i32 %arg0, 0
 ; CHECK-NOT: DIVERGENT
 ; CHECK: ret void

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

-; RUN: opt -mtriple amdgcn-unknown-amdhsa -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 ; RUN: opt -mtriple amdgcn-unknown-amdhsa -passes='print<uniformity>' -disable-output %s 2>&1 | FileCheck %s
 
 define amdgpu_kernel void @hidden_diverge(i32 %n, i32 %a, i32 %b) #0 {
 ; CHECK-LABEL: 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

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

-; RUN: opt -mtriple amdgcn-unknown-amdhsa -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 ; RUN: opt -mtriple amdgcn-unknown-amdhsa -passes='print<uniformity>' -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: 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
 
 H:
   %uni.merge.h = phi i32 [ 0, %entry ], [ %uni.inc, %B ]
   %div.exitx = icmp slt i32 %tid, 0
   br i1 %div.exitx, label %X, label %B ; divergent branch
-; CHECK: DIVERGENT: %div.exitx =  
-; CHECK: DIVERGENT: br i1 %div.exitx, 
+; CHECK: DIVERGENT: %div.exitx =
+; CHECK: DIVERGENT: br i1 %div.exitx,
 
 B:
   %uni.inc = add i32 %uni.merge.h, 1
   %div.exity = icmp sgt i32 %tid, 0
   br i1 %div.exity, label %Y, label %H ; divergent branch
-; CHECK: DIVERGENT: %div.exity =  
-; CHECK: DIVERGENT: br i1 %div.exity, 
+; CHECK: DIVERGENT: %div.exity =
+; CHECK: DIVERGENT: br i1 %div.exity,
 
 X:
   %div.merge.x = phi i32 [ %a, %entry ], [ %uni.merge.h, %H ] ; temporal divergent phi
   br i1 %uni.cond, label %Y, label %exit
 ; CHECK: DIVERGENT: %div.merge.x =
 
 Y:
   %div.merge.y = phi i32 [ 42, %X ], [ %b, %B ]
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
 
 H:
   %uni.merge.h = phi i32 [ 0, %entry ], [ %uni.inc, %B ]
   %div.exitx = icmp slt i32 %tid, 0
   br i1 %div.exitx, label %X, label %B ; divergent branch
-; CHECK: DIVERGENT: %div.exitx =  
-; CHECK: DIVERGENT: br i1 %div.exitx, 
+; CHECK: DIVERGENT: %div.exitx =
+; CHECK: DIVERGENT: br i1 %div.exitx,
 
 B:
   %uni.inc = add i32 %uni.merge.h, 1
   %div.exity = icmp sgt i32 %tid, 0
   br i1 %div.exity, label %Y, label %H ; divergent branch
-; CHECK: DIVERGENT: %div.exity =  
-; CHECK: DIVERGENT: br i1 %div.exity, 
+; CHECK: DIVERGENT: %div.exity =
+; CHECK: DIVERGENT: br i1 %div.exity,
 
 X:
-  %uni.merge.x = phi i32 [ %a, %entry ], [ %b, %H ] 
+  %uni.merge.x = phi i32 [ %a, %entry ], [ %b, %H ]
   br label %exit
 
 Y:
   %uni.merge.y = phi i32 [ %b, %B ]
   br label %exit
 
 exit:
   %div.merge.exit = phi i32 [ %a, %X ], [ %b, %Y ]
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
 
 H:
   %uni.merge.h = phi i32 [ 0, %entry ], [ %uni.inc, %D ]
   br i1 %uni.cond, label %G, label %B
-; CHECK: DIVERGENT: %div.exitx =  
+; CHECK: DIVERGENT: %div.exitx =
 
 B:
-  br i1 %uni.cond, label %X, label %C 
+  br i1 %uni.cond, label %X, label %C
 
 C:
   br i1 %uni.cond, label %Y, label %D
 
 D:
   %uni.inc = add i32 %uni.merge.h, 1
   br label %H
 
 G:
   br i1 %div.exitx, label %C, label %L
-; CHECK: DIVERGENT: br i1 %div.exitx, 
+; CHECK: DIVERGENT: br i1 %div.exitx,
 
 L:
   br i1 %uni.cond, label %D, label %G
 
 X:
   %uni.merge.x = phi i32 [ %a, %entry ], [ %uni.merge.h, %B ]
   br i1 %uni.cond, label %Y, label %exit
 
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
 
 H:
   %uni.merge.h = phi i32 [ 0, %entry ], [ %uni.inc, %C ]
   br i1 %uni.cond, label %G, label %B
-; CHECK: DIVERGENT: %div.exitx =  
+; CHECK: DIVERGENT: %div.exitx =
 
 B:
-  br i1 %uni.cond, label %Y, label %C 
+  br i1 %uni.cond, label %Y, label %C
 
 C:
   %uni.inc = add i32 %uni.merge.h, 1
   br label %H
 
 G:
   br i1 %div.exitx, label %X, label %L ; two-level break
-; CHECK: DIVERGENT: br i1 %div.exitx, 
+; CHECK: DIVERGENT: br i1 %div.exitx,
 
 L:
   br i1 %uni.cond, label %C, label %G
 
 X:
   %div.merge.x = phi i32 [ %a, %entry ], [ %uni.merge.h, %G ] ; temporal divergence
   br label %Y
 ; CHECK: DIVERGENT: %div.merge.x =
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
 
 H:
   %uni.merge.h = phi i32 [ 0, %entry ], [ %uni.inc.d, %D ]
   br i1 %uni.cond, label %G, label %B
-; CHECK: DIVERGENT: %div.exitx =  
+; CHECK: DIVERGENT: %div.exitx =
 
 B:
   %div.merge.b = phi i32 [ 42, %H ], [ %uni.merge.g, %G ]
   br label %D
 ; CHECK: DIVERGENT: %div.merge.b =
 
 G:
   %uni.merge.g = phi i32 [ 123, %H ], [ %uni.inc.l, %L ]
   br i1 %div.exitx, label %B, label %L
-; CHECK: DIVERGENT: br i1 %div.exitx, 
+; CHECK: DIVERGENT: br i1 %div.exitx,
 
 L:
   %uni.inc.l = add i32 %uni.merge.g, 1
   br i1 %uni.cond, label %G, label %D
 
 D:
   %uni.inc.d = add i32 %uni.merge.h, 1
   br i1 %uni.cond, label %X, label %H

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

-; 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
 ; RUN: opt -mtriple amdgcn-unknown-amdhsa -mcpu=tahiti -passes='print<uniformity>' -disable-output %s 2>&1 | FileCheck %s
 ; RUN: opt -mtriple amdgcn-unknown-amdhsa -mcpu=gfx908 -passes='print<uniformity>' -disable-output %s 2>&1 | FileCheck %s
 ; Make sure nothing crashes on targets with or without AGPRs
 
 ; CHECK-LABEL: 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"()

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

-; RUN: opt -mtriple amdgcn-- -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 ; RUN: opt -mtriple amdgcn-- -passes='print<uniformity>' -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/UniformityAnalysis/AMDGPU/intrinsics.ll

-; RUN: opt -mtriple amdgcn-- -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 ; RUN: opt -mtriple amdgcn-- -passes='print<uniformity>' -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(ptr addrspace(1) %out, i32 %src) #0 {
   %swizzle = call i32 @llvm.amdgcn.ds.swizzle(i32 %src, i32 100) #0
   store i32 %swizzle, ptr addrspace(1) %out, align 4
   ret void
 }

llvm/test/Analysis/UniformityAnalysis/AMDGPU/irreducible/branch-outside.ll

 ; RUN: opt %s -mtriple amdgcn-- -passes='print<uniformity>' -disable-output 2>&1 | FileCheck %s
 
 ; CHECK=LABEL: UniformityInfo for function 'basic':
 ; CHECK: CYCLES ASSSUMED DIVERGENT:
 ; CHECK:   depth=1: entries(P T) Q
 define amdgpu_kernel void @basic(i32 %a, i32 %b, i32 %c) {
 entry:
  %cond.uni = icmp slt i32 %a, 0
  %tid = call i32 @llvm.amdgcn.workitem.id.x()
  %cond.div = icmp slt i32 %tid, 0
  br i1 %cond.div, label %T, label %P
 
 P:
 ; CHECK:   DIVERGENT:   %pp.phi =
 ; CHECK: DIVERGENT:      %pp =
   %pp.phi  = phi i32 [ %a, %entry], [ %b, %T ]
   %pp = add i32 %b, 1
   br label %Q
 
 Q:
 ; CHECK: DIVERGENT:   %qq =
 ; CHECK:   DIVERGENT:   %qq.div =
   %qq = add i32 %b, 1
   %qq.div = add i32 %pp.phi, 1
   br i1 %cond.uni, label %T, label %exit
 
 T:
 ; CHECK:   DIVERGENT:   %t.phi =
 ; CHECK: DIVERGENT:     %tt =
   %t.phi = phi i32 [ %qq, %Q ], [ %a, %entry ]
   %tt = add i32 %b, 1
   br label %P
 
 exit:
 ; CHECK-NOT: DIVERGENT:     %ee =
   %ee = add i32 %b, 1
   ret void
 }
 
 ; CHECK=LABEL: UniformityInfo for function 'nested':
 ; CHECK: CYCLES ASSSUMED DIVERGENT:
 ; CHECK:  depth=1: entries(P T) Q A C B
 define amdgpu_kernel void @nested(i32 %a, i32 %b, i32 %c) {
 entry:
  %cond.uni = icmp slt i32 %a, 0
  %tid = call i32 @llvm.amdgcn.workitem.id.x()
  %cond.div = icmp slt i32 %tid, 0
  br i1 %cond.div, label %T, label %P
 
 P:
  %pp.phi  = phi i32 [ %a, %entry], [ %b, %T ]
  %pp = add i32 %b, 1
  br i1 %cond.uni, label %B, label %Q
 
 Q:
   %qq = add i32 %b, 1
   br i1 %cond.uni, label %T, label %exit
 
 A:
   %aa = add i32 %b, 1
   br label %B
 
 B:
   %bb = add i32 %b, 1
   br label %C
 
 C:
   %cc = add i32 %b, 1
   br i1 %cond.uni, label %Q, label %A
 
 T:
   %t.phi = phi i32 [ %qq, %Q ], [ %a, %entry ]
   %tt = add i32 %b, 1
   br i1 %cond.uni, label %A, label %P
 
 exit:
   %ee = add i32 %b, 1
   ret void
 }
 
 declare i32 @llvm.amdgcn.workitem.id.x() #0

llvm/test/Analysis/UniformityAnalysis/AMDGPU/irreducible/diverged-entry-basic.ll

 ; RUN: opt %s -mtriple amdgcn-- -passes='print<uniformity>' -disable-output 2>&1 | FileCheck %s
 
 define amdgpu_kernel void @divergent_cycle_1(i32 %a, i32 %b, i32 %c) {
 ; CHECK-LABEL: UniformityInfo for function 'divergent_cycle_1':
 ; CHECK: CYCLES ASSSUMED DIVERGENT:
 ; CHECK:   depth=1: entries(R P) S Q
 ; CHECK: CYCLES WITH DIVERGENT EXIT:
 ; CHECK:   depth=2: entries(S P) Q
 ; CHECK:   depth=1: entries(R P) S Q
 entry:
  %cond.uni = icmp slt i32 %a, 0
  %tid = call i32 @llvm.amdgcn.workitem.id.x()
  %cond.div = icmp slt i32 %tid, 0
  br i1 %cond.uni, label %P, label %R
 
 P:
 ; CHECK:   DIVERGENT:   %pp.phi =
   %pp.phi  = phi i32 [ %a, %entry], [ %b, %S ]
   %pp = add i32 %b, 1
   br label %Q
 
 Q:
   %qq = add i32 %b, 1
   br i1 %cond.div, label %S, label %R
 
 R:
   %rr = add i32 %b, 1
   br label %S
 
 S:
 ; CHECK:   DIVERGENT:   %s.phi =
   %s.phi = phi i32 [ %qq, %Q ], [ %rr, %R ]
   %ss = add i32 %b, 1
   br i1 %cond.uni, label %exit, label %P
 
 exit:
   %ee = add i32 %b, 1
   ret void
 }
 
 define amdgpu_kernel void @uniform_cycle_1(i32 %a, i32 %b, i32 %c) {
 ; CHECK-LABEL: UniformityInfo for function 'uniform_cycle_1':
 ; CHECK-NOT: CYCLES ASSSUMED DIVERGENT:
 ; CHECK-NOT: CYCLES WITH DIVERGENT EXIT:
 entry:
  %cond.uni = icmp slt i32 %a, 0
  %tid = call i32 @llvm.amdgcn.workitem.id.x()
  %cond.div = icmp slt i32 %tid, 0
  br i1 %cond.uni, label %P, label %T
 
 P:
 ; CHECK-NOT:   DIVERGENT:   %pp.phi = phi i32
   %pp.phi  = phi i32 [ %a, %entry], [ %b, %T ]
   %pp = add i32 %b, 1
   br label %Q
 
 Q:
   %qq = add i32 %b, 1
   br i1 %cond.div, label %S, label %R
 
 R:
   %rr = add i32 %b, 1
   br label %S
 
 S:
 ; CHECK:   DIVERGENT:   %s.phi =
   %s.phi = phi i32 [ %qq, %Q ], [ %rr, %R ]
   %ss = add i32 %b, 1
   br i1 %cond.uni, label %exit, label %T
 
 T:
   %tt = add i32 %b, 1
   br label %P
 
 exit:
   %ee = add i32 %b, 1
   ret void
 }
 
 declare i32 @llvm.amdgcn.workitem.id.x() #0

llvm/test/Analysis/UniformityAnalysis/AMDGPU/irreducible/diverged-entry-headers-nested.ll

 ; RUN: opt %s -mtriple amdgcn-- -passes='print<uniformity>' -disable-output 2>&1 | FileCheck %s
 
 ; These tests have identical control flow graphs with slight changes
 ; that affect cycle-info. There is a minor functional difference in
 ; the branch conditions; but that is not relevant to the tests.
 
 ;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;
 ;;
 ;; The inner cycle has a header (P) that dominates the join, hence
 ;; both cycles are reported as converged.
 ;;
 ;; CHECK-LABEL: UniformityInfo for function 'headers_b_p':
 ;; CHECK-NOT: CYCLES ASSSUMED DIVERGENT:
 ;; CHECK-NOT: CYCLES WITH DIVERGENT EXIT:
 
 define amdgpu_kernel void @headers_b_p(i32 %a, i32 %b, i32 %c) {
 entry:
  %cond.uni = icmp slt i32 %a, 0
  %tid = call i32 @llvm.amdgcn.workitem.id.x()
  %cond.div = icmp slt i32 %tid, 0
  %a.div = add i32 %tid, %a
  br i1 %cond.uni, label %B, label %A
 
 A:
  br label %B
 
 B:
  br i1 %cond.uni, label %C, label %D
 
 C:
  br i1 %cond.uni, label %T, label %P
 
 P:
   %pp.phi  = phi i32 [ %a, %C], [ %b, %T ]
   %pp = add i32 %b, 1
   br i1 %cond.uni, label %R, label %Q
 
 Q:
   %qq = add i32 %b, 1
   br i1 %cond.div, label %S, label %R
 
 R:
   %rr = add i32 %b, 1
   br label %S
 
 S:
   %s.phi = phi i32 [ %qq, %Q ], [ %rr, %R ]
   %ss = add i32 %pp.phi, 1
   br i1 %cond.uni, label %D, label %T
 
 D:
   br i1 %cond.uni, label %exit, label %A
 
 T:
   %tt.phi = phi i32 [ %ss, %S ], [ %a, %C ]
   %tt = add i32 %b, 1
   br label %P
 
 exit:
   %ee = add i32 %b, 1
   ret void
 }
 
 ;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;
 ;;
 ;; Same as previous, but the outer cycle has a different header (A).
 ;; The inner cycle has a header (P) that dominates the join, hence
 ;; both cycles are reported as converged.
 ;;
 ;; CHECK-LABEL: UniformityInfo for function 'headers_a_p':
 ;; CHECK-NOT: CYCLES ASSSUMED DIVERGENT:
 ;; CHECK-NOT: CYCLES WITH DIVERGENT EXIT:
 
 define amdgpu_kernel void @headers_a_p(i32 %a, i32 %b, i32 %c) {
 entry:
  %cond.uni = icmp slt i32 %a, 0
  %tid = call i32 @llvm.amdgcn.workitem.id.x()
  %cond.div = icmp slt i32 %tid, 0
  %a.div = add i32 %tid, %a
  br i1 %cond.uni, label %B, label %A
 
 A:
  br label %B
 
 B:
  br i1 %cond.uni, label %C, label %D
 
 C:
  br i1 %cond.uni, label %T, label %P
 
 P:
   %pp.phi  = phi i32 [ %a, %C], [ %b, %T ]
   %pp = add i32 %b, 1
   br i1 %cond.uni, label %R, label %Q
 
 Q:
   %qq = add i32 %b, 1
   br i1 %cond.div, label %S, label %R
 
 R:
   %rr = add i32 %b, 1
   br label %S
 
 S:
   %s.phi = phi i32 [ %qq, %Q ], [ %rr, %R ]
   %ss = add i32 %pp.phi, 1
   br i1 %cond.uni, label %D, label %T
 
 D:
   br i1 %cond.uni, label %exit, label %A
 
 T:
   %tt.phi = phi i32 [ %ss, %S ], [ %a, %C ]
   %tt = add i32 %b, 1
   br label %P
 
 exit:
   %ee = add i32 %b, 1
   ret void
 }
 
 ;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;
 ;;
 ;; The inner cycle has a header (T) that does not dominate the join.
 ;; The outer cycle has a header (B) that dominates the join. Hence
 ;; only the inner cycle is reported as diverged.
 ;;
 ;; CHECK-LABEL: UniformityInfo for function 'headers_b_t':
 ;; CHECK: CYCLES ASSSUMED DIVERGENT:
 ;; CHECK:   depth=2: entries(T P) S Q R
 ;; CHECK: CYCLES WITH DIVERGENT EXIT:
 ;; CHECK:   depth=1: entries(B A) D T S Q P R C
 
 define amdgpu_kernel void @headers_b_t(i32 %a, i32 %b, i32 %c) {
 entry:
  %cond.uni = icmp slt i32 %a, 0
  %tid = call i32 @llvm.amdgcn.workitem.id.x()
  %cond.div = icmp slt i32 %tid, 0
  %a.div = add i32 %tid, %a
  br i1 %cond.uni, label %A, label %B
 
 A:
  br label %B
 
 B:
  br i1 %cond.uni, label %C, label %D
 
 C:
  br i1 %cond.uni, label %P, label %T
 
 P:
   %pp.phi  = phi i32 [ %a, %C], [ %b, %T ]
   %pp = add i32 %b, 1
   br i1 %cond.uni, label %R, label %Q
 
 Q:
   %qq = add i32 %b, 1
   br i1 %cond.div, label %S, label %R
 
 R:
   %rr = add i32 %b, 1
   br label %S
 
 S:
   %s.phi = phi i32 [ %qq, %Q ], [ %rr, %R ]
   %ss = add i32 %pp.phi, 1
   br i1 %cond.uni, label %D, label %T
 
 D:
   br i1 %cond.uni, label %exit, label %A
 
 T:
   %tt.phi = phi i32 [ %ss, %S ], [ %a, %C ]
   %tt = add i32 %b, 1
   br label %P
 
 exit:
   %ee = add i32 %b, 1
   ret void
 }
 
 ;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;
 ;;
 ;; The cycles have headers (A and T) that do not dominate the join.
 ;; Hence the outermost cycle is reported as diverged.
 ;;
 ;; CHECK-LABEL: UniformityInfo for function 'headers_a_t':
 ;; CHECK: CYCLES ASSSUMED DIVERGENT:
 ;; CHECK:   depth=1: entries(A B) D T S Q P R C
 ;; CHECK-NOT: CYCLES WITH DIVERGENT EXIT:
 
 define amdgpu_kernel void @headers_a_t(i32 %a, i32 %b, i32 %c) {
 entry:
  %cond.uni = icmp slt i32 %a, 0
  %tid = call i32 @llvm.amdgcn.workitem.id.x()
  %cond.div = icmp slt i32 %tid, 0
  %a.div = add i32 %tid, %a
  br i1 %cond.uni, label %B, label %A
 
 A:
  br label %B
 
 B:
  br i1 %cond.uni, label %C, label %D
 
 C:
  br i1 %cond.uni, label %P, label %T
 
 P:
   %pp.phi  = phi i32 [ %a, %C], [ %b, %T ]
   %pp = add i32 %b, 1
   br i1 %cond.uni, label %R, label %Q
 
 Q:
   %qq = add i32 %b, 1
   br i1 %cond.div, label %S, label %R
 
 R:
   %rr = add i32 %b, 1
   br label %S
 
 S:
   %s.phi = phi i32 [ %qq, %Q ], [ %rr, %R ]
   %ss = add i32 %pp.phi, 1
   br i1 %cond.uni, label %D, label %T
 
 D:
   br i1 %cond.uni, label %exit, label %A
 
 T:
   %tt.phi = phi i32 [ %ss, %S ], [ %a, %C ]
   %tt = add i32 %b, 1
   br label %P
 
 exit:
   %ee = add i32 %b, 1
   ret void
 }
 
 declare i32 @llvm.amdgcn.workitem.id.x() #0

llvm/test/Analysis/UniformityAnalysis/AMDGPU/irreducible/diverged-entry-headers.ll

 ; RUN: opt %s -mtriple amdgcn-- -passes='print<uniformity>' -disable-output 2>&1 | FileCheck %s
 
 ; These tests have identical control flow graphs with slight changes
 ; that affect cycle-info. There is a minor functional difference in
 ; the branch conditions; but that is not relevant to the tests.
 
 ;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;
 ;;
 ;; The cycle has a header (T) that does not dominate the join, hence
 ;; the entire cycle is reported as converged.
 ;;
 ;; CHECK-LABEL: UniformityInfo for function 't_header':
 ;; CHECK: CYCLES ASSSUMED DIVERGENT:
 ;; CHECK:   depth=1: entries(T P) S Q R
 
 define amdgpu_kernel void @t_header(i32 %a, i32 %b, i32 %c) {
 entry:
  %cond.uni = icmp slt i32 %a, 0
  %tid = call i32 @llvm.amdgcn.workitem.id.x()
  %cond.div = icmp slt i32 %tid, 0
  %a.div = add i32 %tid, %a
  br i1 %cond.uni, label %P, label %T
 
 P:
 ; CHECK:   DIVERGENT:   %pp.phi =
   %pp.phi  = phi i32 [ %a, %entry], [ %b, %T ]
   %pp = add i32 %b, 1
   br i1 %cond.uni, label %R, label %Q
 
 Q:
   %qq = add i32 %b, 1
   br i1 %cond.div, label %S, label %R
 
 R:
   %rr = add i32 %b, 1
   br label %S
 
 S:
 ; CHECK:   DIVERGENT:   %s.phi =
 ; CHECK:   DIVERGENT:   %ss =
   %s.phi = phi i32 [ %qq, %Q ], [ %rr, %R ]
   %ss = add i32 %pp.phi, 1
   br i1 %cond.uni, label %exit, label %T
 
 T:
 ; CHECK:   DIVERGENT:   %tt.phi =
   %tt.phi = phi i32 [ %ss, %S ], [ %a, %entry ]
   %tt = add i32 %b, 1
   br label %P
 
 exit:
   %ee = add i32 %b, 1
   ret void
 }
 
 ;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;
 ;;
 ;; The cycle has a header (P) that dominates the join, hence
 ;; the cycle is reported as converged.
 ;;
 ;; CHECK-LABEL: UniformityInfo for function 'p_header':
 ;; CHECK-NOT: CYCLES ASSSUMED DIVERGENT:
 
 define amdgpu_kernel void @p_header(i32 %a, i32 %b, i32 %c) {
 entry:
  %cond.uni = icmp slt i32 %a, 0
  %tid = call i32 @llvm.amdgcn.workitem.id.x()
  %cond.div = icmp slt i32 %tid, 0
  br i1 %cond.uni, label %T, label %P
 
 P:
 ; CHECK-NOT:   DIVERGENT:   %pp.phi = phi i32
   %pp.phi  = phi i32 [ %a, %entry], [ %b, %T ]
   %pp = add i32 %b, 1
   br i1 %cond.uni, label %R, label %Q
 
 Q:
   %qq = add i32 %b, 1
   br i1 %cond.div, label %S, label %R
 
 R:
   %rr = add i32 %b, 1
   br label %S
 
 S:
 ; CHECK:   DIVERGENT:   %s.phi =
 ; CHECK-NOT:   DIVERGENT:   %ss = add i32
   %s.phi = phi i32 [ %qq, %Q ], [ %rr, %R ]
   %ss = add i32 %pp.phi, 1
   br i1 %cond.uni, label %exit, label %T
 
 T:
 ; CHECK-NIT:   DIVERGENT:   %tt.phi = phi i32
   %tt.phi = phi i32 [ %ss, %S ], [ %a, %entry ]
   %tt = add i32 %b, 1
   br label %P
 
 exit:
   %ee = add i32 %b, 1
   ret void
 }
 
 declare i32 @llvm.amdgcn.workitem.id.x() #0

llvm/test/Analysis/UniformityAnalysis/AMDGPU/irreducible/exit-divergence.ll

 ; RUN: opt %s -mtriple amdgcn-- -passes='print<uniformity>' -disable-output 2>&1 | FileCheck %s
 
 ; CHECK=LABEL: UniformityInfo for function 'basic':
 ; CHECK-NOT: CYCLES ASSSUMED DIVERGENT:
 ; CHECK: CYCLES WITH DIVERGENT EXIT:
 ; CHECK:   depth=1: entries(P T) Q
 define amdgpu_kernel void @basic(i32 %a, i32 %b, i32 %c) {
 entry:
  %cond.uni = icmp slt i32 %a, 0
  %tid = call i32 @llvm.amdgcn.workitem.id.x()
  %cond.div = icmp slt i32 %tid, 0
  br i1 %cond.uni, label %T, label %P
 
 P:
   %pp.phi.1  = phi i32 [ %a, %entry], [ %b, %T ]
   %pp.phi.2  = phi i32 [ %a, %entry], [ %tt.phi, %T ]
   %pp = add i32 %b, 1
   br label %Q
 
 Q:
   %qq = add i32 %b, 1
   %qq.div.1 = add i32 %pp.phi.2, 1
   %qq.div.2 = add i32 %pp.phi.2, 1
   br i1 %cond.div, label %T, label %exit
 
 T:
   %tt.phi = phi i32 [ %qq, %Q ], [ %a, %entry ]
   %tt = add i32 %b, 1
   br label %P
 
 exit:
 ; CHECK:   DIVERGENT:   %ee.1 =
 ; CHECK:   DIVERGENT:   %xx.2 =
 ; CHECK-NOT: DIVERGENT:     %ee.3 =
   %ee.1 = add i32 %pp.phi.1, 1
   %xx.2 = add i32 %pp.phi.2, 1
   %ee.3 = add i32 %b, 1
   ret void
 }
 
 ; CHECK-LABEL: UniformityInfo for function 'outer_reducible':
 ; CHECK-NOT: CYCLES ASSSUMED DIVERGENT:
 ; CHECK: CYCLES WITH DIVERGENT EXIT:
 ; CHECK:   depth=1: entries(H) P T R Q
 define amdgpu_kernel void @outer_reducible(i32 %a, i32 %b, i32 %c) {
 entry:
  %cond.uni = icmp slt i32 %a, 0
  %tid = call i32 @llvm.amdgcn.workitem.id.x()
  %cond.div = icmp slt i32 %tid, 0
  br label %H
 
 H:
  br i1 %cond.uni, label %T, label %P
 
 P:
   %pp.phi.1  = phi i32 [ %a, %H], [ %b, %T ]
   %pp.phi.2  = phi i32 [ %a, %H], [ %tt.phi, %T ]
   %pp = add i32 %b, 1
   br label %Q
 
 Q:
   %qq = add i32 %b, 1
   %qq.div.1 = add i32 %pp.phi.2, 1
   %qq.div.2 = add i32 %pp.phi.2, 1
   br i1 %cond.div, label %R, label %exit
 
 R:
   br i1 %cond.uni, label %T, label %H
 
 
 T:
   %tt.phi = phi i32 [ %qq, %R ], [ %a, %H ]
   %tt = add i32 %b, 1
   br label %P
 
 exit:
 ; CHECK:   DIVERGENT:   %ee.1 =
 ; CHECK:   DIVERGENT:   %xx.2 =
 ; CHECK-NOT: DIVERGENT:     %ee.3 =
   %ee.1 = add i32 %pp.phi.1, 1
   %xx.2 = add i32 %pp.phi.2, 1
   %ee.3 = add i32 %b, 1
   ret void
 }
 
 ;      entry(div)
 ;      |   \
 ;      H -> B
 ;      ^   /|
 ;      \--C |
 ;          \|
 ;           X
 ;
 ; This has a divergent cycle due to the external divergent branch, but
 ; there are no divergent exits. Hence a use at X is not divergent
 ; unless the def itself is divergent.
 ;
 ; CHECK-LABEL: UniformityInfo for function 'no_divergent_exit':
 ; CHECK: CYCLES ASSSUMED DIVERGENT:
 ; CHECK:   depth=1: entries(H B) C
 ; CHECK-NOT: CYCLES WITH DIVERGENT EXIT:
 define amdgpu_kernel void @no_divergent_exit(i32 %n, i32 %a, i32 %b) #0 {
 entry:
   %tid = call i32 @llvm.amdgcn.workitem.id.x()
   %div.cond = icmp slt i32 %tid, 0
   %uni.cond = icmp slt i32 %a, 0
   br i1 %div.cond, label %B, label %H
 
 H:                                                ; preds = %C, %entry
 ; CHECK: DIVERGENT:  %div.merge.h =
   %div.merge.h = phi i32 [ 0, %entry ], [ %b, %C ]
   br label %B
 
 B:                                                ; preds = %H, %entry
 ; CHECK: DIVERGENT:  %div.merge.b =
   %div.merge.b = phi i32 [ %a, %H ], [ 1, %entry ]
 ; CHECK-NOT: DIVERGENT  %bb =
   %bb = add i32 %a, 1
 ; CHECK-NOT: DIVERGENT:  br i1 %uni.cond, label %X, label %C
   br i1 %uni.cond, label %X, label %C
 
 C:                                                ; preds = %B
 ; CHECK-NOT: DIVERGENT  %cc =
   %cc = add i32 %a, 1
 ; CHECK-NOT: DIVERGENT:  br i1 %uni.cond, label %X, label %H
   br i1 %uni.cond, label %X, label %H
 
 ; CHECK-LABEL: BLOCK X
 X:                                                ; preds = %C, %B
 ; CHECK: DIVERGENT:  %uni.merge.x =
   %uni.merge.x = phi i32 [ %bb, %B ], [%cc, %C ]
 ; CHECK: DIVERGENT: %div.merge.x =
   %div.merge.x = phi i32 [ %div.merge.b, %B ], [%cc, %C ]
   ret void
 }
 
 declare i32 @llvm.amdgcn.workitem.id.x() #0
 
 attributes #0 = { nounwind readnone }

llvm/test/Analysis/UniformityAnalysis/AMDGPU/irreducible/irreducible-1.ll

-; RUN: opt %s -mtriple amdgcn-- -passes='print<divergence>' -disable-output 2>&1 | FileCheck %s
 ; RUN: opt %s -mtriple amdgcn-- -passes='print<uniformity>' -disable-output 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

llvm/test/Analysis/UniformityAnalysis/AMDGPU/irreducible/irreducible-2.ll

 ; RUN: opt %s -mtriple amdgcn-- -passes='print<uniformity>' -disable-output 2>&1 | FileCheck %s
 
 define amdgpu_kernel void @cycle_diverge_enter(i32 %n, i32 %a, i32 %b) #0 {
 ;      entry(div)
 ;      /   \
 ;     H <-> B
 ;           |
 ;           X
 ; CHECK-LABEL: for function 'cycle_diverge_enter':
 ; CHECK-NOT: DIVERGENT: %uni.
 ; CHECK-NOT: DIVERGENT: br i1 %uni.
 
 entry:
   %tid = call i32 @llvm.amdgcn.workitem.id.x()
   %div.cond = icmp slt i32 %tid, 0
   %uni.cond = icmp slt i32 %a, 0
   br i1 %div.cond, label %B, label %H  ; divergent branch
 
 H:
   %div.merge.h = phi i32 [ 0, %entry ], [ %b, %B ]
   br label %B
 ; CHECK: DIVERGENT: %div.merge.h
 
 B:
   %div.merge.b = phi i32 [ %a, %H ], [1, %entry ]
   %div.cond.b = icmp sgt i32 %div.merge.b, 0
   %div.b.inc = add i32 %b, 1
   br i1 %div.cond, label %X, label %H ; divergent branch
 ; CHECK: DIVERGENT: %div.merge.b
 
 X:
   %div.use = add i32 %div.merge.b, 1
   ret void
 ; CHECK: DIVERGENT: %div.use =
 
 }
 
 define amdgpu_kernel void @cycle_diverge_exit(i32 %n, i32 %a, i32 %b) #0 {
 ;      entry
 ;      /   \
 ;     H <-> B(div)
 ;           |
 ;           X
 ;
 ; CHECK-LABEL: for function 'cycle_diverge_exit':
 ; CHECK-NOT: DIVERGENT: %uni.
 ; CHECK-NOT: DIVERGENT: br i1 %uni.
 
 entry:
   %tid = call i32 @llvm.amdgcn.workitem.id.x()
   %div.cond = icmp slt i32 %tid, 0
   %uni.cond = icmp slt i32 %a, 0
   br i1 %uni.cond, label %B, label %H
 
 H:
   %uni.merge.h = phi i32 [ 0, %entry ], [ %b, %B ]
   br label %B
 
 B:
   %uni.merge.b = phi i32 [ %a, %H ], [1, %entry ]
   %uni.cond.b = icmp sgt i32 %uni.merge.b, 0
   %uni.b.inc = add i32 %b, 1
   br i1 %div.cond, label %X, label %H ; divergent branch
 
 X:
   %div.use = add i32 %uni.merge.b, 1
   ret void
 ; CHECK: DIVERGENT: %div.use =
 }
 
 define amdgpu_kernel void @cycle_reentrance(i32 %n, i32 %a, i32 %b) #0 {
 ; For this case, threads enter the cycle from C would take C->D->H,
 ; at the point of H, diverged threads may continue looping in cycle(H-B-D)
 ; until all threads exit the cycle(H-B-D) and cause temporal divergence
 ; exiting at edge H->C. We currently do not analyze such kind of inner
 ; cycle temporal divergence. Instead, we mark all values in the cycle
 ; being divergent conservatively.
 ;      entry--\
 ;       |     |
 ;  ---> H(div)|
 ;  |   / \    /
 ;  |  B   C<--
 ;  ^   \ /
 ;  \----D
 ;       |
 ;       X
 ; CHECK-LABEL: for function 'cycle_reentrance':
 ; CHECK-NOT: DIVERGENT: %uni.
 ; CHECK-NOT: DIVERGENT: br i1 %uni.
 
 entry:
   %tid = call i32 @llvm.amdgcn.workitem.id.x()
   %div.cond = icmp slt i32 %tid, 0
   %uni.cond = icmp slt i32 %a, 0
   br i1 %uni.cond, label %H, label %C
 
 H:
   %div.merge.h = phi i32 [ 0, %entry ], [ %b, %D ]
   br i1 %div.cond, label %B, label %C  ; divergent branch
 
 B:
   %div.inc.b = add i32 %div.merge.h, 1
 ; CHECK: DIVERGENT: %div.inc.b
   br label %D
 
 C:
   %div.merge.c = phi i32 [0, %entry], [%div.merge.h, %H]
   %div.inc.c = add i32 %div.merge.c, 2
 ; CHECK: DIVERGENT: %div.inc.c
   br label %D
 
 D:
   %div.merge.d = phi i32 [ %div.inc.b, %B ], [ %div.inc.c, %C ]
 ; CHECK: DIVERGENT: %div.merge.d
   br i1 %uni.cond, label %X, label %H
 
 X:
   ret void
 }
 
 define amdgpu_kernel void @cycle_reentrance2(i32 %n, i32 %a, i32 %b) #0 {
 ; This is mostly the same as cycle_reentrance, the only difference is
 ; the successor order, thus different dfs visiting order. This is just
 ; make sure we are doing uniform analysis correctly under different dfs
 ; order.
 ;      entry--\
 ;       |     |
 ;  ---> H(div)|
 ;  |   / \    /
 ;  |  B   C<--
 ;  ^   \ /
 ;  \----D
 ;       |
 ;       X
 ; CHECK-LABEL: for function 'cycle_reentrance2':
 ; CHECK-NOT: DIVERGENT: %uni.
 ; CHECK-NOT: DIVERGENT: br i1 %uni.
 
 entry:
   %tid = call i32 @llvm.amdgcn.workitem.id.x()
   %div.cond = icmp slt i32 %tid, 0
   %uni.cond = icmp slt i32 %a, 0
   br i1 %uni.cond, label %C, label %H
 
 H:
   %div.merge.h = phi i32 [ 0, %entry ], [ %b, %D ]
   br i1 %div.cond, label %B, label %C  ; divergent branch
 
 B:
   %div.inc.b = add i32 %div.merge.h, 1
 ; CHECK: DIVERGENT: %div.inc.b
   br label %D
 
 C:
   %div.merge.c = phi i32 [0, %entry], [%div.merge.h, %H]
   %div.inc.c = add i32 %div.merge.c, 2
 ; CHECK: DIVERGENT: %div.inc.c
   br label %D
 
 D:
   %div.merge.d = phi i32 [ %div.inc.b, %B ], [ %div.inc.c, %C ]
 ; CHECK: DIVERGENT: %div.merge.d
   br i1 %uni.cond, label %X, label %H
 
 X:
   ret void
 }
 
 define amdgpu_kernel void @cycle_join_dominated_by_diverge(i32 %n, i32 %a, i32 %b) #0 {
 ; the join-node D is dominated by diverge point H2
 ;      entry
 ;       | |
 ;  --> H1 |
 ;  |     \|
 ;  |      H2(div)
 ;  |     / \
 ;  |    B   C
 ;  ^     \ /
 ;  \------D
 ;         |
 ;         X
 ; CHECK-LABEL: for function 'cycle_join_dominated_by_diverge':
 ; CHECK-NOT: DIVERGENT: %uni.
 ; CHECK-NOT: DIVERGENT: br i1 %uni.
 
 entry:
   %tid = call i32 @llvm.amdgcn.workitem.id.x()
   %div.cond = icmp slt i32 %tid, 0
   %uni.cond = icmp slt i32 %a, 0
   br i1 %uni.cond, label %H1, label %H2
 
 H1:
   %uni.merge.h1 = phi i32 [ 0, %entry ], [ %b, %D ]
   br label %H2
 
 H2:
   %uni.merge.h2 = phi i32 [ 0, %entry ], [ %b, %H1 ]
   br i1 %div.cond, label %B, label %C  ; divergent branch
 
 B:
   %uni.inc.b = add i32 %uni.merge.h2, 1
   br label %D
 
 C:
   %uni.inc.c = add i32 %uni.merge.h2, 2
   br label %D
 
 D:
   %div.merge.d = phi i32 [ %uni.inc.b, %B ], [ %uni.inc.c, %C ]
 ; CHECK: DIVERGENT: %div.merge.d
   br i1 %uni.cond, label %X, label %H1
 
 X:
   ret void
 }
 
 define amdgpu_kernel void @cycle_join_dominated_by_entry(i32 %n, i32 %a, i32 %b) #0 {
 ; the join-node D is dominated by cycle entry H2
 ;      entry
 ;       | |
 ;  --> H1 |
 ;  |     \|
 ;  |      H2 -----
 ;  |      |      |
 ;  |      A(div) |
 ;  |     / \     v
 ;  |    B   C   /
 ;  ^     \ /   /
 ;  \------D <-/
 ;         |
 ;         X
 ; CHECK-LABEL: for function 'cycle_join_dominated_by_entry':
 ; CHECK-NOT: DIVERGENT: %uni.
 ; CHECK-NOT: DIVERGENT: br i1 %uni.
 
 entry:
   %tid = call i32 @llvm.amdgcn.workitem.id.x()
   %div.cond = icmp slt i32 %tid, 0
   %uni.cond = icmp slt i32 %a, 0
   br i1 %uni.cond, label %H1, label %H2
 
 H1:
   %uni.merge.h1 = phi i32 [ 0, %entry ], [ %b, %D ]
   br label %H2
 
 H2:
   %uni.merge.h2 = phi i32 [ 0, %entry ], [ %b, %H1 ]
   br i1 %uni.cond, label %A, label %D
 
 A:
   br i1 %div.cond, label %B, label %C  ; divergent branch
 
 B:
   %uni.inc.b = add i32 %uni.merge.h2, 1
   br label %D
 
 C:
   %uni.inc.c = add i32 %uni.merge.h2, 2
   br label %D
 
 D:
   %div.merge.d = phi i32 [ %uni.inc.b, %B ], [ %uni.inc.c, %C ], [%uni.merge.h2, %H2]
 ; CHECK: DIVERGENT: %div.merge.d
   br i1 %uni.cond, label %X, label %H1
 
 X:
   ret void
 }
 
 define amdgpu_kernel void @cycle_join_not_dominated(i32 %n, i32 %a, i32 %b) #0 {
 ; if H is the header, the sync label propagation may stop at join node D.
 ; But join node D is not dominated by divergence starting block C, and also
 ; not dominated by any entries(H/C). So we conservatively mark all the values
 ; in the cycle divergent for now.
 ;      entry
 ;       |  |
 ;  ---> H  |
 ;  |   / \ v
 ;  |  B<--C(div)
 ;  ^   \ /
 ;  \----D
 ;       |
 ;       X
 ; CHECK-LABEL: for function 'cycle_join_not_dominated':
 ; CHECK-NOT: DIVERGENT: %uni.
 
 entry:
   %tid = call i32 @llvm.amdgcn.workitem.id.x()
   %div.cond = icmp slt i32 %tid, 0
   %uni.cond = icmp slt i32 %a, 0
   br i1 %uni.cond, label %C, label %H
 
 H:
   %div.merge.h = phi i32 [ 0, %entry ], [ %b, %D ]
   br i1 %uni.cond, label %B, label %C
 
 B:
   %div.merge.b = phi i32 [ 0, %H ], [ %b, %C ]
   %div.inc.b = add i32 %div.merge.b, 1
 ; CHECK: DIVERGENT: %div.inc.b
   br label %D
 
 C:
   %div.merge.c = phi i32 [0, %entry], [%div.merge.h, %H]
   %div.inc.c = add i32 %div.merge.c, 2
 ; CHECK: DIVERGENT: %div.inc.c
   br i1 %div.cond, label %D, label %B  ; divergent branch
 
 D:
   %div.merge.d = phi i32 [ %div.inc.b, %B ], [ %div.inc.c, %C ]
 ; CHECK: DIVERGENT: %div.merge.d
   br i1 %uni.cond, label %X, label %H
 
 X:
   ret void
 }
 
 define amdgpu_kernel void @cycle_join_not_dominated2(i32 %n, i32 %a, i32 %b) #0 {
 ; This is mostly the same as cycle_join_not_dominated, the only difference is
 ; the dfs visiting order, so the cycle analysis result is different.
 ;      entry
 ;       |  |
 ;  ---> H  |
 ;  |   / \ v
 ;  |  B<--C(div)
 ;  ^   \ /
 ;  \----D
 ;       |
 ;       X
 ; CHECK-LABEL: for function 'cycle_join_not_dominated2':
 ; CHECK-NOT: DIVERGENT: %uni.
 
 entry:
   %tid = call i32 @llvm.amdgcn.workitem.id.x()
   %div.cond = icmp slt i32 %tid, 0
   %uni.cond = icmp slt i32 %a, 0
   br i1 %uni.cond, label %H, label %C
 
 H:
   %div.merge.h = phi i32 [ 0, %entry ], [ %b, %D ]
   br i1 %uni.cond, label %B, label %C
 
 B:
   %div.merge.b = phi i32 [ 0, %H ], [ %b, %C ]
   %div.inc.b = add i32 %div.merge.b, 1
 ; CHECK: DIVERGENT: %div.inc.b
   br label %D
 
 C:
   %div.merge.c = phi i32 [0, %entry], [%div.merge.h, %H]
   %div.inc.c = add i32 %div.merge.c, 2
 ; CHECK: DIVERGENT: %div.inc.c
   br i1 %div.cond, label %D, label %B  ; divergent branch
 
 D:
   %div.merge.d = phi i32 [ %div.inc.b, %B ], [ %div.inc.c, %C ]
 ; CHECK: DIVERGENT: %div.merge.d
   br i1 %uni.cond, label %X, label %H
 
 X:
   ret void
 }
 
 define amdgpu_kernel void @natural_loop_two_backedges(i32 %n, i32 %a, i32 %b) #0 {
 ; FIXME: the uni.merge.h can be viewed as uniform.
 ; CHECK-LABEL: for function 'natural_loop_two_backedges':
 
 entry:
   %tid = call i32 @llvm.amdgcn.workitem.id.x()
   %div.cond = icmp slt i32 %tid, 0
   %uni.cond = icmp slt i32 %a, 0
   br label %H
 
 H:
   %uni.merge.h = phi i32 [ 0, %entry ], [ %uni.inc, %B ], [ %uni.inc, %C]
   %uni.inc = add i32 %uni.merge.h, 1
   br label %B
 
 B:
   br i1 %div.cond, label %C, label %H
 
 C:
   br i1 %uni.cond, label %X, label %H
 
 X:
   ret void
 }
 
 define amdgpu_kernel void @natural_loop_two_backedges2(i32 %n, i32 %a, i32 %b) #0 {
 ; FIXME: the uni.merge.h can be viewed as uniform.
 ; CHECK-LABEL: for function 'natural_loop_two_backedges2':
 
 entry:
   %tid = call i32 @llvm.amdgcn.workitem.id.x()
   %div.cond = icmp slt i32 %tid, 0
   %uni.cond = icmp slt i32 %a, 0
   br label %H
 
 H:
   %uni.merge.h = phi i32 [ 0, %entry ], [ %uni.inc, %B ], [ %uni.inc, %C]
   %uni.inc = add i32 %uni.merge.h, 1
   br i1 %uni.cond, label %B, label %D
 
 B:
   br i1 %div.cond, label %C, label %H
 
 C:
   br label %H
 
 D:
   br i1 %uni.cond, label %B, label %X
 
 X:
   ret void
 }
 
 define amdgpu_kernel void @cycle_enter_nested(i32 %n, i32 %a, i32 %b) #0 {
 ;
 ;   entry(div)
 ;       |   \
 ;   --> H1   |
 ;  /    |    |
 ;  | -> H2   |
 ;  | |  |    /
 ;  | \--B <--
 ;  ^    |
 ;  \----C
 ;       |
 ;       X
 ; CHECK-LABEL: for function 'cycle_enter_nested':
 ; CHECK-NOT: DIVERGENT: %uni.
 ; CHECK-NOT: DIVERGENT: br i1 %uni.
 
 entry:
   %tid = call i32 @llvm.amdgcn.workitem.id.x()
   %div.cond = icmp slt i32 %tid, 0
   %uni.cond = icmp slt i32 %a, 0
   br i1 %div.cond, label %B, label %H1
 
 H1:
   %div.merge.h1 = phi i32 [ 1, %entry ], [ %b, %C ]
   br label %H2
 ; CHECK: DIVERGENT: %div.merge.h1
 
 H2:
   %div.merge.h2 = phi i32 [ 2, %B ], [ %a, %H1 ]
 ; CHECK: DIVERGENT: %div.merge.h2
   br label %B
 
 B:
   %div.merge.b = phi i32 [0, %entry], [%a, %H2]
 ; CHECK: DIVERGENT: %div.merge.b
   br i1 %uni.cond, label %C, label %H2
 
 C:
   br i1 %uni.cond, label %X, label %H1
 
 X:
   ret void
 }
 
 define amdgpu_kernel void @cycle_inner_exit_enter(i32 %n, i32 %a, i32 %b) #0 {
 ;          entry
 ;        /      \
 ;       E1-> A-> E2
 ;       ^    |    \
 ;       |    E3-> E4
 ;       |    ^   /
 ;       |     \ /
 ;       C <----B
 ;       |
 ;       X
 ; CHECK-LABEL: for function 'cycle_inner_exit_enter':
 ; CHECK-NOT: DIVERGENT: %uni.
 ; CHECK-NOT: DIVERGENT: br i1 %uni.
 
 entry:
   %tid = call i32 @llvm.amdgcn.workitem.id.x()
   %div.cond = icmp slt i32 %tid, 0
   %uni.cond = icmp slt i32 %a, 0
   br i1 %uni.cond, label %E2, label %E1
 
 E1:
   %div.merge.e1 = phi i32 [ 1, %entry ], [ %b, %C ]
   br label %A
 ; CHECK: DIVERGENT: %div.merge.e1
 
 A:
   br i1 %uni.cond, label %E2, label %E3
 
 E2:
   %div.merge.e2 = phi i32 [ 2, %entry ], [ %a, %A ]
 ; CHECK: DIVERGENT: %div.merge.e2
   br label %E4
 
 E3:
   %div.merge.e3 = phi i32 [ 0, %A ], [ %b, %B ]
 ; CHECK: DIVERGENT: %div.merge.e3
   br label %E4
 
 E4:
   %div.merge.e4 = phi i32 [ 0, %E2 ], [ %a, %E3 ]
 ; CHECK: DIVERGENT: %div.merge.e4
   br label %B
 
 B:
   br i1 %div.cond, label %C, label %E3
 
 C:
   br i1 %uni.cond, label %X, label %E1
 
 X:
   ret void
 }
 
 define amdgpu_kernel void @cycle_inner_exit_enter2(i32 %n, i32 %a, i32 %b) #0 {
 ; This case is almost the same as cycle_inner_exit_enter, with only different
 ; dfs visiting order, thus different cycle hierarchy.
 ;          entry
 ;        /      \
 ;       E1-> A-> E2
 ;       ^    |    \
 ;       |    E3-> E4
 ;       |    ^   /
 ;       |     \ /
 ;       C <----B
 ;       |
 ;       X
 ; CHECK-LABEL: for function 'cycle_inner_exit_enter2':
 ; CHECK-NOT: DIVERGENT: %uni.
 ; CHECK-NOT: DIVERGENT: br i1 %uni.
 
 entry:
   %tid = call i32 @llvm.amdgcn.workitem.id.x()
   %div.cond = icmp slt i32 %tid, 0
   %uni.cond = icmp slt i32 %a, 0
   br i1 %uni.cond, label %E1, label %E2
 
 E1:
   %div.merge.e1 = phi i32 [ 1, %entry ], [ %b, %C ]
   br label %A
 ; CHECK: DIVERGENT: %div.merge.e1
 
 A:
   br i1 %uni.cond, label %E2, label %E3
 
 E2:
   %div.merge.e2 = phi i32 [ 2, %entry ], [ %a, %A ]
 ; CHECK: DIVERGENT: %div.merge.e2
   br label %E4
 
 E3:
   %div.merge.e3 = phi i32 [ 0, %A ], [ %b, %B ]
 ; CHECK: DIVERGENT: %div.merge.e3
   br label %E4
 
 E4:
   %div.merge.e4 = phi i32 [ 0, %E2 ], [ %a, %E3 ]
 ; CHECK: DIVERGENT: %div.merge.e4
   br label %B
 
 B:
   br i1 %div.cond, label %C, label %E3
 
 C:
   br i1 %uni.cond, label %X, label %E1
 
 X:
   ret void
 }
 
 define amdgpu_kernel void @always_uniform() {
 ; CHECK-LABEL: UniformityInfo for function 'always_uniform':
 ; CHECK: CYCLES ASSSUMED DIVERGENT:
 ; CHECK:   depth=1: entries(bb2 bb3)
 
 bb:
   %inst = tail call i32 @llvm.amdgcn.mbcnt.hi(i32 0, i32 0)
   %inst1 = icmp ugt i32 %inst, 0
   br i1 %inst1, label %bb3, label %bb2
 ; CHECK:   DIVERGENT:   %inst = tail call i32 @llvm.amdgcn.mbcnt.hi(i32 0, i32 0)
 ; CHECK:   DIVERGENT:   %inst1 = icmp ugt i32 %inst, 0
 ; CHECK:   DIVERGENT:   br i1 %inst1, label %bb3, label %bb2
 
 bb2:                                              ; preds = %bb3, %bb
   br label %bb3
 
 bb3:                                              ; preds = %bb2, %bb
   %inst4 = tail call i64 @llvm.amdgcn.icmp.i64.i16(i16 0, i16 0, i32 0)
   %inst5 = trunc i64 %inst4 to i32
   %inst6 = and i32 0, %inst5
   br label %bb2
 ; CHECK-LABEL: BLOCK bb3
 ; CHECK-NOT: DIVERGENT: {{.*}} call i64 @llvm.amdgcn.icmp.i64.i16
 ; CHECK:   DIVERGENT:   %inst5 = trunc i64 %inst4 to i32
 ; CHECK:   DIVERGENT:   %inst6 = and i32 0, %inst5
 }
 
 declare i32 @llvm.amdgcn.mbcnt.hi(i32, i32)
 
 declare i64 @llvm.amdgcn.icmp.i64.i16(i16, i16, i32 immarg)
 
 declare i32 @llvm.amdgcn.workitem.id.x() #0
 
 attributes #0 = { nounwind readnone }

llvm/test/Analysis/UniformityAnalysis/AMDGPU/irreducible/reducible-headers.ll

 ; RUN: opt %s -mtriple amdgcn-- -passes='print<uniformity>' -disable-output 2>&1 | FileCheck %s
 
 ;
 ;                         Entry
 ;                           |
 ;                           v
 ;                  -------->H---------
 ;                  |        |        |
 ;                  |        v        |
 ;                  |    --->T----    |
 ;                  |    |       |    |
 ;                  |    |       V    |
 ;                  S<---R       P <---
 ;                  ^    ^       |
 ;                  |    |  Div  |
 ;                  |    --- Q <--
 ;                  |        |
 ;                  |        v
 ;                  -------- U
 ;                           |
 ;                           v
 ;                          Exit
 ;
 ; The divergent branch is at Q that exits an irreducible cycle with
 ; entries T and P nested inside a reducible cycle with header H. R is
 ; assigned label R, which reaches P. S is a join node with label S. If
 ; this is propagated to P via H, then P is incorrectly recognized as a
 ; join, making the inner cycle divergent. P is always executed
 ; convergently -- either by threads that reconverged at header H, or
 ; by threads that are still executing the inner cycle. Thus, any PHI
 ; at P should not be marked divergent.
 
 define amdgpu_kernel void @nested_irreducible(i32 %a, i32 %b, i32 %c) {
 ; CHECK=LABEL: UniformityInfo for function 'nested_irreducible':
 ; CHECK-NOT: CYCLES ASSSUMED DIVERGENT:
 ; CHECK: CYCLES WITH DIVERGENT EXIT:
 ; CHECK-DAG:   depth=2: entries(P T) R Q
 ; CHECK-DAG:   depth=1: entries(H) S P T R Q U
 entry:
   %cond.uni = icmp slt i32 %a, 0
   %tid = call i32 @llvm.amdgcn.workitem.id.x()
   %cond.div = icmp slt i32 %tid, 0
   br label %H
 
 H:
  br i1 %cond.uni, label %T, label %P
 
 P:
 ; CHECK-LABEL: BLOCK P
 ; CHECK-NOT:   DIVERGENT:   %pp.phi =
 ; CHECK-NOT: DIVERGENT:      %pp =
   %pp.phi  = phi i32 [ %a, %H], [ %b, %T ]
   %pp = add i32 %b, 1
   br label %Q
 
 Q:
 ; CHECK-LABEL: BLOCK Q
 ; CHECK-NOT: DIVERGENT:   %qq =
 ; CHECK-NOT:   DIVERGENT:   %qq.uni =
   %qq = add i32 %b, 1
   %qq.uni = add i32 %pp.phi, 1
   br i1 %cond.div, label %R, label %U
 
 R:
   br i1 %cond.uni, label %S, label %T
 
 T:
 ; CHECK-LABEL: BLOCK T
 ; CHECK-NOT:   DIVERGENT:   %tt.phi =
 ; CHECK-NOT: DIVERGENT:     %tt =
   %tt.phi = phi i32 [ %qq, %R ], [ %a, %H ]
   %tt = add i32 %b, 1
   br label %P
 
 S:
 ; CHECK-LABEL: BLOCK S
 ; CHECK:   DIVERGENT:   %ss.phi =
 ; CHECK-NOT: DIVERGENT:     %ss =
   %ss.phi = phi i32 [ %qq.uni, %U ], [ %a, %R ]
   %ss = add i32 %b, 1
   br label %H
 
 U:
   br i1 %cond.uni, label %S, label %exit
 
 exit:
 ; CHECK: DIVERGENT:     %ee.div =
 ; CHECK-NOT: DIVERGENT:     %ee =
   %ee.div =  add i32 %qq.uni, 1
   %ee = add i32 %b, 1
   ret void
 }
 
 ;
 ;                         Entry
 ;                           |
 ;                           v
 ;               -->-------->H---------
 ;               |  ^        |        |
 ;               |  |        |        |
 ;               |  |        |        |
 ;               |  |        |        |
 ;               |  |        v        V
 ;               |  R<-------T-->U--->P
 ;               |          Div       |
 ;               |                    |
 ;               ----------- Q <-------
 ;                           |
 ;                           v
 ;                          Exit
 ;
 ; This is a reducible cycle with a divergent branch at T. Disjoint
 ; paths eventually join at the header H, which is assigned label H.
 ; Node P is assigned label U. If the header label were propagated to
 ; P, it will be incorrectly recgonized as a join. P is always executed
 ; convergently -- either by threads that reconverged at header H, or
 ; by threads that diverged at T (and eventually reconverged at H).
 ; Thus, any PHI at P should not be marked divergent.
 
 define amdgpu_kernel void @header_label_1(i32 %a, i32 %b, i32 %c) {
 ; CHECK=LABEL: UniformityInfo for function 'header_label_1':
 ; CHECK-NOT: CYCLES ASSSUMED DIVERGENT:
 ; CHECK: CYCLES WITH DIVERGENT EXIT:
 ; CHECK:  depth=1: entries(H) Q P U T R
 entry:
   %cond.uni = icmp slt i32 %a, 0
   %tid = call i32 @llvm.amdgcn.workitem.id.x()
   %cond.div = icmp slt i32 %tid, 0
   br label %H
 
 H:
   br i1 %cond.uni, label %T, label %P
 
 P:
 ; CHECK-LABEL: BLOCK P
 ; CHECK-NOT:   DIVERGENT:   %pp.phi =
 ; CHECK-NOT: DIVERGENT:      %pp =
   %pp.phi  = phi i32 [ %a, %H], [ %b, %U ]
   %pp = add i32 %b, 1
   br label %Q
 
 Q:
 ; CHECK-LABEL: BLOCK Q
 ; CHECK-NOT: DIVERGENT:   %qq =
 ; CHECK-NOT:   DIVERGENT:   %qq.uni =
   %qq = add i32 %b, 1
   %qq.uni = add i32 %pp.phi, 1
   br i1 %cond.uni, label %exit, label %H
 
 R:
   br label %H
 
 T:
   br i1 %cond.div, label %R, label %U
 
 U:
   br label %P
 
 exit:
 ; CHECK-LABEL: BLOCK exit
 ; CHECK: DIVERGENT:     %ee.div =
 ; CHECK-NOT: DIVERGENT:     %ee =
   %ee.div =  add i32 %qq.uni, 1
   %ee = add i32 %b, 1
   ret void
 }
 
 ;        entry
 ;            |
 ;        --> H1
 ;        |   | \
 ;        |   | H2(div)
 ;        |   \ / \
 ;        |    B   C
 ;        ^     \ /
 ;        \------D
 ;               |
 ;               X
 ;
 ; This is a reducible cycle with a divergent branch at H2. Disjoint
 ; paths eventually join at the header D, which is assigned label D.
 ; Node B is assigned label B. If the header label D were propagated to
 ; B, it will be incorrectly recgonized as a join. B is always executed
 ; convergently -- either by threads that reconverged at header H1, or
 ; by threads that diverge at H2 (and eventually reconverged at H1).
 ; Thus, any PHI at B should not be marked divergent.
 
 define amdgpu_kernel void @header_label_2(i32 %a, i32 %b, i32 %c) {
 ; CHECK-LABEL: UniformityInfo for function 'header_label_2':
 ; CHECK-NOT: CYCLES ASSSUMED DIVERGENT:
 ; CHECK-NOT: CYCLES WITH DIVERGENT EXIT:
 entry:
   %cond.uni = icmp slt i32 %a, 0
   %tid = call i32 @llvm.amdgcn.workitem.id.x()
   %cond.div = icmp slt i32 %tid, 0
   br label %H1
 
 H1:
   br i1 %cond.uni, label %B, label %H2
 
 H2:
   br i1 %cond.div, label %B, label %C
 
 B:
 ; CHECK-LABEL: BLOCK B
 ; CHECK-NOT: DIVERGENT:     %bb.phi =
   %bb.phi = phi i32 [%a, %H1], [%b, %H2]
   br label %D
 
 C:
   br label %D
 
 D:
 ; CHECK-LABEL: BLOCK D
 ; CHECK: DIVERGENT:     %dd.phi =
   %dd.phi = phi i32 [%a, %B], [%b, %C]
   br i1 %cond.uni, label %exit, label %H1
 
 exit:
   %ee.1 = add i32 %dd.phi, 1
   %ee.2 = add i32 %b, 1
   ret void
 }
 
 declare i32 @llvm.amdgcn.workitem.id.x() #0

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

-; RUN: opt -mtriple amdgcn-unknown-amdhsa -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 ; RUN: opt -mtriple amdgcn-unknown-amdhsa -passes='print<uniformity>' -disable-output %s 2>&1 | FileCheck %s
 
 ; 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/UniformityAnalysis/AMDGPU/join-at-loop-heart.ll

-; RUN: opt -mtriple amdgcn-unknown-amdhsa -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 ; RUN: opt -mtriple amdgcn-unknown-amdhsa -passes='print<uniformity>' -disable-output %s 2>&1 | FileCheck %s
 
 ; CHECK: DIVERGENT:  %phi.h = phi i32 [ 0, %entry ], [ %inc, %C ], [ %inc, %D ], [ %inc, %E ]
 ; CHECK: DIVERGENT:       %tid = call i32 @llvm.amdgcn.workitem.id.x()
 ; CHECK: DIVERGENT:       %div.cond = icmp slt i32 %tid, 0
 ; CHECK: DIVERGENT:  %inc = add i32 %phi.h, 1
 ; CHECK: DIVERGENT:       br i1 %div.cond, label %C, label %D
 

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

-; RUN: opt -mtriple amdgcn-- -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 ; RUN: opt -mtriple amdgcn-- -passes='print<uniformity>' -disable-output %s 2>&1 | FileCheck %s
 
 ; CHECK-LABEL: for function 'test_amdgpu_ps':
 ; CHECK-DAG: DIVERGENT:  ptr addrspace(4) %arg0
 ; CHECK-DAG: DIVERGENT:  <2 x i32> %arg3
 ; CHECK-DAG: DIVERGENT:  <3 x i32> %arg4
 ; CHECK-DAG: DIVERGENT:  float %arg5
 ; CHECK-DAG: DIVERGENT:  i32 %arg6

llvm/test/Analysis/UniformityAnalysis/AMDGPU/lit.local.cfg

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

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

-; RUN: opt -mtriple amdgcn-mesa-mesa3d -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 ; RUN: opt -mtriple amdgcn-mesa-mesa3d -passes='print<uniformity>' -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/UniformityAnalysis/AMDGPU/llvm.amdgcn.image.atomic.ll

-; RUN: opt -mtriple amdgcn-mesa-mesa3d -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 ; RUN: opt -mtriple amdgcn-mesa-mesa3d -passes='print<uniformity>' -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/UniformityAnalysis/AMDGPU/no-return-blocks.ll

-; RUN: opt -mtriple amdgcn-- -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 ; RUN: opt -mtriple amdgcn-- -passes='print<uniformity>' -disable-output %s 2>&1 | FileCheck %s
 
 ; CHECK: DIVERGENT:  %tmp5 = getelementptr inbounds float, ptr addrspace(1) %arg, i64 %tmp2
 ; CHECK: DIVERGENT:  %tmp10 = load volatile float, ptr addrspace(1) %tmp5, align 4
 ; CHECK: DIVERGENT:  %tmp11 = load volatile float, ptr 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(ptr addrspace(1) noalias nocapture readonly %arg, ptr addrspace(1) noalias nocapture readonly %arg1) #0 {

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

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

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

-; RUN: opt -mtriple amdgcn-unknown-amdhsa -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 ; RUN: opt -mtriple amdgcn-unknown-amdhsa -passes='print<uniformity>' -disable-output %s 2>&1 | FileCheck %s
 
 ; CHECK: DIVERGENT:       %.126.i355.i = phi i1 [ false, %bb5 ], [ true, %bb4 ]
 
 ; Function Attrs: nounwind readnone speculatable
 declare i32 @llvm.amdgcn.workitem.id.x() #0
 
 define protected amdgpu_kernel void @_Z23krnl_GPUITSFitterKerneli() {

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

-; RUN: opt -mtriple amdgcn-unknown-amdhsa -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 ; RUN: opt -mtriple amdgcn-unknown-amdhsa -passes='print<uniformity>' -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: 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, 
+; 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 {

 G:
   br label %H
 
 H:
   %uni.merge.h = phi i32 [ 0, %G ], [ %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, 
+; CHECK: DIVERGENT: %div.exitx =
+; CHECK: DIVERGENT: br i1 %div.exitx,
 
 X:
   %div.user = add i32 %uni.inc, 5
   br i1 %uni.cond, label %G, label %Y
 
 Y:
   %div.alsouser = add i32 %uni.inc, 5
   ret void

 H:
   %uni.merge.h = phi i32 [ 0, %G ], [ %uni.inc, %H ]
   %uni.inc = add i32 %uni.merge.h, 1
   br i1 %uni.cond, label %X, label %H ; divergent branch
 
 X:
   %uni.user = add i32 %uni.inc, 5
   %div.exity = icmp slt i32 %tid, 0
-; CHECK: DIVERGENT: %div.exity =  
+; CHECK: DIVERGENT: %div.exity =
   br i1 %div.exity, label %G, label %Y
-; CHECK: DIVERGENT: br i1 %div.exity, 
+; CHECK: DIVERGENT: br i1 %div.exity,
 
 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: 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
 
 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, 
+; CHECK: DIVERGENT: %div.exitx =
+; CHECK: DIVERGENT: br i1 %div.exitx,
 
 X:
   br label %G
 
 G:
   %div.user = add i32 %uni.inc, 5
   br i1 %uni.cond, label %G, label %Y
 
entry:
   %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, 
+; CHECK: DIVERGENT: %div.exitx =
+; CHECK: DIVERGENT: br i1 %div.exitx,
 
 X:
   br label %G
 
 G:
   %div.user = add i32 %uni.inc, 5
   br i1 %uni.cond, label %G, label %Y
 
 Y:
   ret void
 }
 
 
 declare i32 @llvm.amdgcn.workitem.id.x() #0
 
 attributes #0 = { nounwind readnone }

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

-; RUN: opt -mtriple amdgcn-unknown-amdhsa -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 ; RUN: opt -mtriple amdgcn-unknown-amdhsa -passes='print<uniformity>' -disable-output %s 2>&1 | FileCheck %s
 
 ; 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/UniformityAnalysis/AMDGPU/unreachable-loop-block.ll

-; RUN: opt -mtriple amdgcn-- -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 ; RUN: opt -mtriple amdgcn-- -passes='print<uniformity>' -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

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

-; RUN: opt -mtriple amdgcn-unknown-amdhsa -passes='print<divergence>' -disable-output %s 2>&1 | FileCheck %s
 ; RUN: opt -mtriple amdgcn-unknown-amdhsa -passes='print<uniformity>' -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
 

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

-; RUN: opt %s -passes='print<divergence>' -disable-output 2>&1 | FileCheck %s
 ; RUN: opt %s -passes='print<uniformity>' -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: 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: %cond =
 ; CHECK: DIVERGENT: br i1 %cond,
 A:
   %defAtA = add i32 %n, 1 ; uniform
-; CHECK-NOT: DIVERGENT: %defAtA = 
+; CHECK-NOT: DIVERGENT: %defAtA =
   br label %C
 B:
   %defAtB = add i32 %n, 2 ; uniform
-; CHECK-NOT: DIVERGENT: %defAtB = 
+; CHECK-NOT: DIVERGENT: %defAtB =
   br label %C
 C:
   %defAtC = phi i32 [ %defAtA, %A ], [ %defAtB, %B ] ; divergent
 ; CHECK: DIVERGENT: %defAtC =
   br label %D
 
 D:
   %i = phi i32 [0, %C], [ %i.inc, %E ] ; uniform

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

-; RUN: opt %s -passes='print<divergence>' -disable-output 2>&1 | FileCheck %s
 ; RUN: opt %s -passes='print<uniformity>' -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: for function 'no_diverge'

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

-; RUN: opt %s -passes='print<divergence>' -disable-output 2>&1 | FileCheck %s
 ; RUN: opt %s -passes='print<uniformity>' -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: for function 'hidden_diverge'
 entry:

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

-; RUN: opt %s -passes='print<divergence>' -disable-output 2>&1 | FileCheck %s
 ; RUN: opt %s -passes='print<uniformity>' -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.

llvm/test/Analysis/UniformityAnalysis/NVPTX/lit.local.cfg

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

llvm/test/CodeGen/AMDGPU/always-uniform.ll

-; RUN: opt -mtriple amdgcn-amdhsa -mcpu=gfx90a -passes=legacy-divergence-analysis < %s -S 2>&1 | FileCheck -check-prefix=OPT %s
 ; RUN: llc -mtriple amdgcn-amdhsa -mcpu=fiji -amdgpu-scalarize-global-loads -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
 
 declare i32 @llvm.amdgcn.workitem.id.x()
 declare i32 @llvm.amdgcn.readfirstlane(i32)
 
-; OPT-LABEL: define amdgpu_kernel void @readfirstlane_uniform(
-; OPT-NEXT:    %tid = tail call i32 @llvm.amdgcn.workitem.id.x()
-; OPT-NEXT:    %scalar = tail call i32 @llvm.amdgcn.readfirstlane(i32 %tid)
-; OPT-NEXT:    %idx = zext i32 %scalar to i64
-; OPT-NEXT:    %gep0 = getelementptr inbounds float, ptr addrspace(1) %0, i64 %idx
-; OPT-NEXT:    %val = load float, ptr addrspace(1) %gep0, align 4
-; OPT-NEXT:    %gep1 = getelementptr inbounds float, ptr addrspace(1) %1, i64 10
-; OPT-NEXT:    store float %val, ptr addrspace(1) %gep1, align 4
-; OPT-NEXT:    ret void
-;
 ; GCN-LABEL: readfirstlane_uniform
 ; GCN:  s_load_dwordx4 s[[[IN_ADDR:[0-9]+]]:3], s[4:5], 0x0
 ; GCN:  v_readfirstlane_b32 s[[SCALAR:[0-9]+]], v0
 ; GCN:  s_add_u32 s[[LOAD_ADDR:[0-9]+]], s[[IN_ADDR]], s[[SCALAR]]
 ; GCN:  s_load_dword s{{[0-9]+}}, s[[[LOAD_ADDR]]
 
 define amdgpu_kernel void @readfirstlane_uniform(ptr addrspace(1) noalias nocapture readonly, ptr addrspace(1) noalias nocapture readonly) {
   %tid = tail call i32 @llvm.amdgcn.workitem.id.x()
   %scalar = tail call i32 @llvm.amdgcn.readfirstlane(i32 %tid)
   %idx = zext i32 %scalar to i64
   %gep0 = getelementptr inbounds float, ptr addrspace(1) %0, i64 %idx
   %val = load float, ptr addrspace(1) %gep0, align 4
   %gep1 = getelementptr inbounds float, ptr addrspace(1) %1, i64 10
   store float %val, ptr addrspace(1) %gep1, align 4
   ret void
 }

llvm/test/CodeGen/AMDGPU/smrd.ll

loop:
   br i1 %cc, label %exit, label %loop
 
 exit:
   ret float %sum.next
 }
 
 
 ; GCN-LABEL: {{^}}smrd_uniform_loop2:
-; (this test differs from smrd_uniform_loop by the more complex structure of phis,
-; which used to confuse the DivergenceAnalysis after structurization)
+; (this test differs from smrd_uniform_loop by the more complex structure of phis)
 ;
 ; TODO: we should keep the loop counter in an SGPR and use an S_BUFFER_LOAD
 ;
 ; GCN: buffer_load_dword
 define amdgpu_ps float @smrd_uniform_loop2(<4 x i32> inreg %desc, i32 %bound, i32 %bound.a) #0 {
 main_body:
   br label %loop
 

llvm/unittests/Analysis/CMakeLists.txt

set(ANALYSIS_TEST_SOURCES
   BlockFrequencyInfoTest.cpp
   BranchProbabilityInfoTest.cpp
   CallGraphTest.cpp
   CaptureTrackingTest.cpp
   CFGTest.cpp
   CGSCCPassManagerTest.cpp
   ConstraintSystemTest.cpp
   DDGTest.cpp
-  DivergenceAnalysisTest.cpp
   DomTreeUpdaterTest.cpp
   GlobalsModRefTest.cpp
   FunctionPropertiesAnalysisTest.cpp
   InlineCostTest.cpp
   IRSimilarityIdentifierTest.cpp
   IVDescriptorsTest.cpp
   LazyCallGraphTest.cpp
   LoadsTest.cpp

llvm/unittests/Analysis/DivergenceAnalysisTest.cpp

-//===- DivergenceAnalysisTest.cpp - DivergenceAnalysis unit tests ---------===//
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-//===----------------------------------------------------------------------===//
-
-#include "llvm/ADT/SmallVector.h"
-#include "llvm/Analysis/AssumptionCache.h"
-#include "llvm/Analysis/DivergenceAnalysis.h"
-#include "llvm/Analysis/LoopInfo.h"
-#include "llvm/Analysis/PostDominators.h"
-#include "llvm/Analysis/SyncDependenceAnalysis.h"
-#include "llvm/Analysis/TargetLibraryInfo.h"
-#include "llvm/AsmParser/Parser.h"
-#include "llvm/IR/Constants.h"
-#include "llvm/IR/Dominators.h"
-#include "llvm/IR/GlobalVariable.h"
-#include "llvm/IR/IRBuilder.h"
-#include "llvm/IR/InstIterator.h"
-#include "llvm/IR/LLVMContext.h"
-#include "llvm/IR/LegacyPassManager.h"
-#include "llvm/IR/Module.h"
-#include "llvm/IR/Verifier.h"
-#include "llvm/Support/SourceMgr.h"
-#include "gtest/gtest.h"
-
-namespace llvm {
-namespace {
-
-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 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) {}
-
-  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 DivergenceAnalysisImpl(F, nullptr, *DT, *LI, *SDA, IsLCSSA);
-  }
-
-  void runWithDA(
-      Module &M, StringRef FuncName, bool IsLCSSA,
-      function_ref<void(Function &F, LoopInfo &LI, DivergenceAnalysisImpl &DA)>
-          Test) {
-    auto *F = M.getFunction(FuncName);
-    ASSERT_NE(F, nullptr) << "Could not find " << FuncName;
-    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);
-
-  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
-  DA.compute();
-  EXPECT_FALSE(DA.hasDetectedDivergence());
-
-  // Detected divergence after marking
-  Argument &arg = *F->arg_begin();
-  DA.markDivergent(arg);
-
-  EXPECT_TRUE(DA.hasDetectedDivergence());
-  EXPECT_TRUE(DA.isDivergent(arg));
-
-  DA.compute();
-  EXPECT_TRUE(DA.hasDetectedDivergence());
-  EXPECT_TRUE(DA.isDivergent(arg));
-}
-
-TEST_F(DivergenceAnalysisTest, DANoLCSSA) {
-  LLVMContext C;
-  SMDiagnostic Err;
-
-  std::unique_ptr<Module> M = parseAssemblyString(
-      "target datalayout = \"e-m:e-p:32:32-f64:32:64-f80:32-n8:16:32-S128\" "
-      " "
-      "define i32 @f_1(i8* nocapture %arr, i32 %n, i32* %A, i32* %B) "
-      "    local_unnamed_addr { "
-      "entry: "
-      "  br label %loop.ph "
-      " "
-      "loop.ph: "
-      "  br label %loop "
-      " "
-      "loop: "
-      "  %iv0 = phi i32 [ %iv0.inc, %loop ], [ 0, %loop.ph ] "
-      "  %iv1 = phi i32 [ %iv1.inc, %loop ], [ -2147483648, %loop.ph ] "
-      "  %iv0.inc = add i32 %iv0, 1 "
-      "  %iv1.inc = add i32 %iv1, 3 "
-      "  %cond.cont = icmp slt i32 %iv0, %n "
-      "  br i1 %cond.cont, label %loop, label %for.end.loopexit "
-      " "
-      "for.end.loopexit: "
-      "  ret i32 %iv0 "
-      "} ",
-      Err, C);
-
-  Function *F = M->getFunction("f_1");
-  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);
-
-  DA.compute();
-  EXPECT_TRUE(DA.hasDetectedDivergence());
-
-  // Verify that "ret %iv.0" is divergent
-  auto ItBlock = F->begin();
-  std::advance(ItBlock, 3);
-  auto &ExitBlock = *GetBlockByName("for.end.loopexit", *F);
-  auto &RetInst = *cast<ReturnInst>(ExitBlock.begin());
-  EXPECT_TRUE(DA.isDivergent(RetInst));
-}
-
-TEST_F(DivergenceAnalysisTest, DALCSSA) {
-  LLVMContext C;
-  SMDiagnostic Err;
-
-  std::unique_ptr<Module> M = parseAssemblyString(
-      "target datalayout = \"e-m:e-p:32:32-f64:32:64-f80:32-n8:16:32-S128\" "
-      " "
-      "define i32 @f_lcssa(i8* nocapture %arr, i32 %n, i32* %A, i32* %B) "
-      "    local_unnamed_addr { "
-      "entry: "
-      "  br label %loop.ph "
-      " "
-      "loop.ph: "
-      "  br label %loop "
-      " "
-      "loop: "
-      "  %iv0 = phi i32 [ %iv0.inc, %loop ], [ 0, %loop.ph ] "
-      "  %iv1 = phi i32 [ %iv1.inc, %loop ], [ -2147483648, %loop.ph ] "
-      "  %iv0.inc = add i32 %iv0, 1 "
-      "  %iv1.inc = add i32 %iv1, 3 "
-      "  %cond.cont = icmp slt i32 %iv0, %n "
-      "  br i1 %cond.cont, label %loop, label %for.end.loopexit "
-      " "
-      "for.end.loopexit: "
-      "  %val.ret = phi i32 [ %iv0, %loop ] "
-      "  br label %detached.return "
-      " "
-      "detached.return: "
-      "  ret i32 %val.ret "
-      "} ",
-      Err, C);
-
-  Function *F = M->getFunction("f_lcssa");
-  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);
-
-  DA.compute();
-  EXPECT_TRUE(DA.hasDetectedDivergence());
-
-  // Verify that "ret %iv.0" is divergent
-  auto ItBlock = F->begin();
-  std::advance(ItBlock, 4);
-  auto &ExitBlock = *GetBlockByName("detached.return", *F);
-  auto &RetInst = *cast<ReturnInst>(ExitBlock.begin());
-  EXPECT_TRUE(DA.isDivergent(RetInst));
-}
-
-TEST_F(DivergenceAnalysisTest, DAJoinDivergence) {
-  LLVMContext C;
-  SMDiagnostic Err;
-
-  std::unique_ptr<Module> M = parseAssemblyString(
-      "target datalayout = \"e-m:e-p:32:32-f64:32:64-f80:32-n8:16:32-S128\" "
-      " "
-      "define void @f_1(i1 %a, i1 %b, i1 %c) "
-      "    local_unnamed_addr { "
-      "A: "
-      "  br i1 %a, label %B, label %C "
-      " "
-      "B: "
-      "  br i1 %b, label %C, label %D "
-      " "
-      "C: "
-      "  %c.join = phi i32 [ 0, %A ], [ 1, %B ] "
-      "  br i1 %c, label %D, label %E "
-      " "
-      "D: "
-      "  %d.join = phi i32 [ 0, %B ], [ 1, %C ] "
-      "  br label %E "
-      " "
-      "E: "
-      "  %e.join = phi i32 [ 0, %C ], [ 1, %D ] "
-      "  ret void "
-      "} "
-      " "
-      "define void @f_2(i1 %a, i1 %b, i1 %c) "
-      "    local_unnamed_addr { "
-      "A: "
-      "  br i1 %a, label %B, label %E "
-      " "
-      "B: "
-      "  br i1 %b, label %C, label %D "
-      " "
-      "C: "
-      "  br label %D "
-      " "
-      "D: "
-      "  %d.join = phi i32 [ 0, %B ], [ 1, %C ] "
-      "  br label %E "
-      " "
-      "E: "
-      "  %e.join = phi i32 [ 0, %A ], [ 1, %D ] "
-      "  ret void "
-      "} "
-      " "
-      "define void @f_3(i1 %a, i1 %b, i1 %c)"
-      "    local_unnamed_addr { "
-      "A: "
-      "  br i1 %a, label %B, label %C "
-      " "
-      "B: "
-      "  br label %C "
-      " "
-      "C: "
-      "  %c.join = phi i32 [ 0, %A ], [ 1, %B ] "
-      "  br i1 %c, label %D, label %E "
-      " "
-      "D: "
-      "  br label %E "
-      " "
-      "E: "
-      "  %e.join = phi i32 [ 0, %C ], [ 1, %D ] "
-      "  ret void "
-      "} ",
-      Err, C);
-
-  // Maps divergent conditions to the basic blocks whose Phi nodes become
-  // divergent. Blocks need to be listed in IR order.
-  using SmallBlockVec = SmallVector<const BasicBlock *, 4>;
-  using InducedDivJoinMap = std::map<const Value *, SmallBlockVec>;
-
-  // Actual function performing the checks.
-  auto CheckDivergenceFunc = [this](Function &F,
-                                    InducedDivJoinMap &ExpectedDivJoins) {
-    for (auto &ItCase : ExpectedDivJoins) {
-      auto *DivVal = ItCase.first;
-      auto DA = buildDA(F, false);
-      DA.markDivergent(*DivVal);
-      DA.compute();
-
-      // List of basic blocks that shall host divergent Phi nodes.
-      auto ItDivJoins = ItCase.second.begin();
-
-      for (auto &BB : F) {
-        auto *Phi = dyn_cast<PHINode>(BB.begin());
-        if (!Phi)
-          continue;
-
-        if (ItDivJoins != ItCase.second.end() && &BB == *ItDivJoins) {
-          EXPECT_TRUE(DA.isDivergent(*Phi));
-          // Advance to next block with expected divergent PHI node.
-          ++ItDivJoins;
-        } else {
-          EXPECT_FALSE(DA.isDivergent(*Phi));
-        }
-      }
-    }
-  };
-
-  {
-    auto *F = M->getFunction("f_1");
-    auto ItBlocks = F->begin();
-    ItBlocks++; // Skip A
-    ItBlocks++; // Skip B
-    auto *C = &*ItBlocks++;
-    auto *D = &*ItBlocks++;
-    auto *E = &*ItBlocks;
-
-    auto ItArg = F->arg_begin();
-    auto *AArg = &*ItArg++;
-    auto *BArg = &*ItArg++;
-    auto *CArg = &*ItArg;
-
-    InducedDivJoinMap DivJoins;
-    DivJoins.emplace(AArg, SmallBlockVec({C, D, E}));
-    DivJoins.emplace(BArg, SmallBlockVec({D, E}));
-    DivJoins.emplace(CArg, SmallBlockVec({E}));
-
-    CheckDivergenceFunc(*F, DivJoins);
-  }
-
-  {
-    auto *F = M->getFunction("f_2");
-    auto ItBlocks = F->begin();
-    ItBlocks++; // Skip A
-    ItBlocks++; // Skip B
-    ItBlocks++; // Skip C
-    auto *D = &*ItBlocks++;
-    auto *E = &*ItBlocks;
-
-    auto ItArg = F->arg_begin();
-    auto *AArg = &*ItArg++;
-    auto *BArg = &*ItArg++;
-    auto *CArg = &*ItArg;
-
-    InducedDivJoinMap DivJoins;
-    DivJoins.emplace(AArg, SmallBlockVec({E}));
-    DivJoins.emplace(BArg, SmallBlockVec({D}));
-    DivJoins.emplace(CArg, SmallBlockVec({}));
-
-    CheckDivergenceFunc(*F, DivJoins);
-  }
-
-  {
-    auto *F = M->getFunction("f_3");
-    auto ItBlocks = F->begin();
-    ItBlocks++; // Skip A
-    ItBlocks++; // Skip B
-    auto *C = &*ItBlocks++;
-    ItBlocks++; // Skip D
-    auto *E = &*ItBlocks;
-
-    auto ItArg = F->arg_begin();
-    auto *AArg = &*ItArg++;
-    auto *BArg = &*ItArg++;
-    auto *CArg = &*ItArg;
-
-    InducedDivJoinMap DivJoins;
-    DivJoins.emplace(AArg, SmallBlockVec({C}));
-    DivJoins.emplace(BArg, SmallBlockVec({}));
-    DivJoins.emplace(CArg, SmallBlockVec({E}));
-
-    CheckDivergenceFunc(*F, DivJoins);
-  }
-}
-
-TEST_F(DivergenceAnalysisTest, DASwitchUnreachableDefault) {
-  LLVMContext C;
-  SMDiagnostic Err;
-
-  std::unique_ptr<Module> M = parseAssemblyString(
-      "target datalayout = \"e-m:e-p:32:32-f64:32:64-f80:32-n8:16:32-S128\" "
-      " "
-      "define void @switch_unreachable_default(i32 %cond) local_unnamed_addr { "
-      "entry: "
-      "  switch i32 %cond, label %sw.default [ "
-      "    i32 0, label %sw.bb0 "
-      "    i32 1, label %sw.bb1 "
-      "  ] "
-      " "
-      "sw.bb0: "
-      "  br label %sw.epilog "
-      " "
-      "sw.bb1: "
-      "  br label %sw.epilog "
-      " "
-      "sw.default: "
-      "  unreachable "
-      " "
-      "sw.epilog: "
-      "  %div.dbl = phi double [ 0.0, %sw.bb0], [ -1.0, %sw.bb1 ] "
-      "  ret void "
-      "}",
-      Err, C);
-
-  auto *F = M->getFunction("switch_unreachable_default");
-  auto &CondArg = *F->arg_begin();
-  auto DA = buildDA(*F, false);
-
-  EXPECT_FALSE(DA.hasDetectedDivergence());
-
-  DA.markDivergent(CondArg);
-  DA.compute();
-
-  // Still %CondArg is divergent.
-  EXPECT_TRUE(DA.hasDetectedDivergence());
-
-  // The join uni.dbl is not divergent (see D52221)
-  auto &ExitBlock = *GetBlockByName("sw.epilog", *F);
-  auto &DivDblPhi = *cast<PHINode>(ExitBlock.begin());
-  EXPECT_TRUE(DA.isDivergent(DivDblPhi));
-}
-
-} // end anonymous namespace
-} // end namespace llvm

llvm/utils/gn/secondary/llvm/lib/Analysis/BUILD.gn

sources = [
     "CycleAnalysis.cpp",
     "DDG.cpp",
     "DDGPrinter.cpp",
     "Delinearization.cpp",
     "DemandedBits.cpp",
     "DependenceAnalysis.cpp",
     "DependenceGraphBuilder.cpp",
     "DevelopmentModeInlineAdvisor.cpp",
-    "DivergenceAnalysis.cpp",
     "DomPrinter.cpp",
     "DomTreeUpdater.cpp",
     "DominanceFrontier.cpp",
     "FunctionPropertiesAnalysis.cpp",
     "GlobalsModRef.cpp",
     "GuardUtils.cpp",
     "HeatUtils.cpp",
     "IRSimilarityIdentifier.cpp",
sources = [
     "InstructionSimplify.cpp",
     "InteractiveModelRunner.cpp",
     "Interval.cpp",
     "IntervalPartition.cpp",
     "LazyBlockFrequencyInfo.cpp",
     "LazyBranchProbabilityInfo.cpp",
     "LazyCallGraph.cpp",
     "LazyValueInfo.cpp",
-    "LegacyDivergenceAnalysis.cpp",
     "Lint.cpp",
     "Loads.cpp",
     "Local.cpp",
     "LoopAccessAnalysis.cpp",
     "LoopAnalysisManager.cpp",
     "LoopCacheAnalysis.cpp",
     "LoopInfo.cpp",
     "LoopNestAnalysis.cpp",

llvm/utils/gn/secondary/llvm/unittests/Analysis/BUILD.gn

sources = [
     "BlockFrequencyInfoTest.cpp",
     "BranchProbabilityInfoTest.cpp",
     "CFGTest.cpp",
     "CGSCCPassManagerTest.cpp",
     "CallGraphTest.cpp",
     "CaptureTrackingTest.cpp",
     "ConstraintSystemTest.cpp",
     "DDGTest.cpp",
-    "DivergenceAnalysisTest.cpp",
     "DomTreeUpdaterTest.cpp",
     "FunctionPropertiesAnalysisTest.cpp",
     "GlobalsModRefTest.cpp",
     "IRSimilarityIdentifierTest.cpp",
     "IVDescriptorsTest.cpp",
     "InlineCostTest.cpp",
     "LazyCallGraphTest.cpp",
     "LoadsTest.cpp",