From 35617ed4cb81ca442a519fd4134ff28b90f06655 Mon Sep 17 00:00:00 2001 From: Nicolai Haehnle Date: Thu, 30 Aug 2018 14:21:36 +0000 Subject: [PATCH] [NFC] Rename the DivergenceAnalysis to LegacyDivergenceAnalysis Summary: This is patch 1 of the new DivergenceAnalysis (https://reviews.llvm.org/D50433). The purpose of this patch is to free up the name DivergenceAnalysis for the new generic implementation. The generic implementation class will be shared by specialized divergence analysis classes. Patch by: Simon Moll Reviewed By: nhaehnle Subscribers: jvesely, jholewinski, arsenm, nhaehnle, mgorny, jfb, llvm-commits Differential Revision: https://reviews.llvm.org/D50434 Change-Id: Ie8146b11be2c50d5312f30e11c7a3036a15b48cb llvm-svn: 341071 --- ...vergenceAnalysis.h => LegacyDivergenceAnalysis.h} | 18 +++++++++--------- llvm/include/llvm/Analysis/Passes.h | 4 ++-- llvm/include/llvm/Analysis/TargetTransformInfo.h | 2 +- llvm/include/llvm/CodeGen/SelectionDAG.h | 6 +++--- llvm/include/llvm/CodeGen/TargetLowering.h | 4 ++-- llvm/include/llvm/InitializePasses.h | 2 +- llvm/include/llvm/LinkAllPasses.h | 2 +- llvm/lib/Analysis/Analysis.cpp | 2 +- llvm/lib/Analysis/CMakeLists.txt | 2 +- ...enceAnalysis.cpp => LegacyDivergenceAnalysis.cpp} | 20 ++++++++++---------- llvm/lib/CodeGen/SelectionDAG/SelectionDAG.cpp | 2 +- llvm/lib/CodeGen/SelectionDAG/SelectionDAGISel.cpp | 2 +- .../Target/AMDGPU/AMDGPUAnnotateUniformValues.cpp | 10 +++++----- llvm/lib/Target/AMDGPU/AMDGPUCodeGenPrepare.cpp | 10 +++++----- llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp | 6 +++--- .../lib/Target/AMDGPU/AMDGPULowerKernelArguments.cpp | 1 - .../Target/AMDGPU/AMDGPUUnifyDivergentExitNodes.cpp | 12 ++++++------ llvm/lib/Target/AMDGPU/SIAnnotateControlFlow.cpp | 10 +++++----- llvm/lib/Target/AMDGPU/SIISelLowering.cpp | 4 ++-- llvm/lib/Target/AMDGPU/SIISelLowering.h | 2 +- llvm/lib/Transforms/Scalar/LoopUnswitch.cpp | 8 ++++---- llvm/lib/Transforms/Scalar/StructurizeCFG.cpp | 12 ++++++------ .../AMDGPU/atomics.ll | 0 .../AMDGPU/intrinsics.ll | 0 .../AMDGPU/kernel-args.ll | 6 +++--- .../AMDGPU/lit.local.cfg | 0 .../AMDGPU/llvm.amdgcn.buffer.atomic.ll | 0 .../AMDGPU/llvm.amdgcn.image.atomic.ll | 0 .../AMDGPU/no-return-blocks.ll | 0 .../AMDGPU/phi-undef.ll | 0 .../AMDGPU/unreachable-loop-block.ll | 0 .../AMDGPU/workitem-intrinsics.ll | 0 .../NVPTX/diverge.ll | 14 +++++++------- .../NVPTX/lit.local.cfg | 0 34 files changed, 80 insertions(+), 81 deletions(-) rename llvm/include/llvm/Analysis/{DivergenceAnalysis.h => LegacyDivergenceAnalysis.h} (72%) rename llvm/lib/Analysis/{DivergenceAnalysis.cpp => LegacyDivergenceAnalysis.cpp} (94%) rename llvm/test/Analysis/{DivergenceAnalysis => LegacyDivergenceAnalysis}/AMDGPU/atomics.ll (100%) rename llvm/test/Analysis/{DivergenceAnalysis => LegacyDivergenceAnalysis}/AMDGPU/intrinsics.ll (100%) rename llvm/test/Analysis/{DivergenceAnalysis => LegacyDivergenceAnalysis}/AMDGPU/kernel-args.ll (80%) rename llvm/test/Analysis/{DivergenceAnalysis => LegacyDivergenceAnalysis}/AMDGPU/lit.local.cfg (100%) rename llvm/test/Analysis/{DivergenceAnalysis => LegacyDivergenceAnalysis}/AMDGPU/llvm.amdgcn.buffer.atomic.ll (100%) rename llvm/test/Analysis/{DivergenceAnalysis => LegacyDivergenceAnalysis}/AMDGPU/llvm.amdgcn.image.atomic.ll (100%) rename llvm/test/Analysis/{DivergenceAnalysis => LegacyDivergenceAnalysis}/AMDGPU/no-return-blocks.ll (100%) rename llvm/test/Analysis/{DivergenceAnalysis => LegacyDivergenceAnalysis}/AMDGPU/phi-undef.ll (100%) rename llvm/test/Analysis/{DivergenceAnalysis => LegacyDivergenceAnalysis}/AMDGPU/unreachable-loop-block.ll (100%) rename llvm/test/Analysis/{DivergenceAnalysis => LegacyDivergenceAnalysis}/AMDGPU/workitem-intrinsics.ll (100%) rename llvm/test/Analysis/{DivergenceAnalysis => LegacyDivergenceAnalysis}/NVPTX/diverge.ll (90%) rename llvm/test/Analysis/{DivergenceAnalysis => LegacyDivergenceAnalysis}/NVPTX/lit.local.cfg (100%) diff --git a/llvm/include/llvm/Analysis/DivergenceAnalysis.h b/llvm/include/llvm/Analysis/LegacyDivergenceAnalysis.h similarity index 72% rename from llvm/include/llvm/Analysis/DivergenceAnalysis.h rename to llvm/include/llvm/Analysis/LegacyDivergenceAnalysis.h index 328c8645..173ca28 100644 --- a/llvm/include/llvm/Analysis/DivergenceAnalysis.h +++ b/llvm/include/llvm/Analysis/LegacyDivergenceAnalysis.h @@ -1,4 +1,4 @@ -//===- llvm/Analysis/DivergenceAnalysis.h - Divergence Analysis -*- C++ -*-===// +//===- llvm/Analysis/LegacyDivergenceAnalysis.h - KernelDivergence Analysis -*- C++ -*-===// // // The LLVM Compiler Infrastructure // @@ -7,14 +7,14 @@ // //===----------------------------------------------------------------------===// // -// The divergence analysis is an LLVM pass which can be used to find out -// if a branch instruction in a GPU program is divergent or not. It can help +// 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_DIVERGENCE_ANALYSIS_H -#define LLVM_ANALYSIS_DIVERGENCE_ANALYSIS_H +#ifndef LLVM_ANALYSIS_LEGACY_DIVERGENCE_ANALYSIS_H +#define LLVM_ANALYSIS_LEGACY_DIVERGENCE_ANALYSIS_H #include "llvm/ADT/DenseSet.h" #include "llvm/IR/Function.h" @@ -22,12 +22,12 @@ namespace llvm { class Value; -class DivergenceAnalysis : public FunctionPass { +class LegacyDivergenceAnalysis : public FunctionPass { public: static char ID; - DivergenceAnalysis() : FunctionPass(ID) { - initializeDivergenceAnalysisPass(*PassRegistry::getPassRegistry()); + LegacyDivergenceAnalysis() : FunctionPass(ID) { + initializeLegacyDivergenceAnalysisPass(*PassRegistry::getPassRegistry()); } void getAnalysisUsage(AnalysisUsage &AU) const override; @@ -58,4 +58,4 @@ private: }; } // End llvm namespace -#endif //LLVM_ANALYSIS_DIVERGENCE_ANALYSIS_H \ No newline at end of file +#endif //LLVM_ANALYSIS_LEGACY_DIVERGENCE_ANALYSIS_H diff --git a/llvm/include/llvm/Analysis/Passes.h b/llvm/include/llvm/Analysis/Passes.h index 09b28a0..081dd50 100644 --- a/llvm/include/llvm/Analysis/Passes.h +++ b/llvm/include/llvm/Analysis/Passes.h @@ -61,10 +61,10 @@ namespace llvm { //===--------------------------------------------------------------------===// // - // createDivergenceAnalysisPass - This pass determines which branches in a GPU + // createLegacyDivergenceAnalysisPass - This pass determines which branches in a GPU // program are divergent. // - FunctionPass *createDivergenceAnalysisPass(); + FunctionPass *createLegacyDivergenceAnalysisPass(); //===--------------------------------------------------------------------===// // diff --git a/llvm/include/llvm/Analysis/TargetTransformInfo.h b/llvm/include/llvm/Analysis/TargetTransformInfo.h index 59657cc..9cf6ff9 100644 --- a/llvm/include/llvm/Analysis/TargetTransformInfo.h +++ b/llvm/include/llvm/Analysis/TargetTransformInfo.h @@ -289,7 +289,7 @@ public: /// Returns whether V is a source of divergence. /// /// This function provides the target-dependent information for - /// the target-independent DivergenceAnalysis. DivergenceAnalysis first + /// the target-independent LegacyDivergenceAnalysis. LegacyDivergenceAnalysis first /// builds the dependency graph, and then runs the reachability algorithm /// starting with the sources of divergence. bool isSourceOfDivergence(const Value *V) const; diff --git a/llvm/include/llvm/CodeGen/SelectionDAG.h b/llvm/include/llvm/CodeGen/SelectionDAG.h index a118c58..dd486837 100644 --- a/llvm/include/llvm/CodeGen/SelectionDAG.h +++ b/llvm/include/llvm/CodeGen/SelectionDAG.h @@ -28,7 +28,7 @@ #include "llvm/ADT/iterator.h" #include "llvm/ADT/iterator_range.h" #include "llvm/Analysis/AliasAnalysis.h" -#include "llvm/Analysis/DivergenceAnalysis.h" +#include "llvm/Analysis/LegacyDivergenceAnalysis.h" #include "llvm/CodeGen/DAGCombine.h" #include "llvm/CodeGen/FunctionLoweringInfo.h" #include "llvm/CodeGen/ISDOpcodes.h" @@ -229,7 +229,7 @@ class SelectionDAG { LLVMContext *Context; CodeGenOpt::Level OptLevel; - DivergenceAnalysis * DA = nullptr; + LegacyDivergenceAnalysis * DA = nullptr; FunctionLoweringInfo * FLI = nullptr; /// The function-level optimization remark emitter. Used to emit remarks @@ -382,7 +382,7 @@ public: /// Prepare this SelectionDAG to process code in the given MachineFunction. void init(MachineFunction &NewMF, OptimizationRemarkEmitter &NewORE, Pass *PassPtr, const TargetLibraryInfo *LibraryInfo, - DivergenceAnalysis * Divergence); + LegacyDivergenceAnalysis * Divergence); void setFunctionLoweringInfo(FunctionLoweringInfo * FuncInfo) { FLI = FuncInfo; diff --git a/llvm/include/llvm/CodeGen/TargetLowering.h b/llvm/include/llvm/CodeGen/TargetLowering.h index 2f5ac04..e49cd3f 100644 --- a/llvm/include/llvm/CodeGen/TargetLowering.h +++ b/llvm/include/llvm/CodeGen/TargetLowering.h @@ -29,7 +29,7 @@ #include "llvm/ADT/STLExtras.h" #include "llvm/ADT/SmallVector.h" #include "llvm/ADT/StringRef.h" -#include "llvm/Analysis/DivergenceAnalysis.h" +#include "llvm/Analysis/LegacyDivergenceAnalysis.h" #include "llvm/CodeGen/DAGCombine.h" #include "llvm/CodeGen/ISDOpcodes.h" #include "llvm/CodeGen/RuntimeLibcalls.h" @@ -2655,7 +2655,7 @@ public: virtual bool isSDNodeSourceOfDivergence(const SDNode *N, FunctionLoweringInfo *FLI, - DivergenceAnalysis *DA) const { + LegacyDivergenceAnalysis *DA) const { return false; } diff --git a/llvm/include/llvm/InitializePasses.h b/llvm/include/llvm/InitializePasses.h index d67b1d4..9e1694d 100644 --- a/llvm/include/llvm/InitializePasses.h +++ b/llvm/include/llvm/InitializePasses.h @@ -119,7 +119,6 @@ void initializeDependenceAnalysisPass(PassRegistry&); void initializeDependenceAnalysisWrapperPassPass(PassRegistry&); void initializeDetectDeadLanesPass(PassRegistry&); void initializeDivRemPairsLegacyPassPass(PassRegistry&); -void initializeDivergenceAnalysisPass(PassRegistry&); void initializeDomOnlyPrinterPass(PassRegistry&); void initializeDomOnlyViewerPass(PassRegistry&); void initializeDomPrinterPass(PassRegistry&); @@ -191,6 +190,7 @@ void initializeLazyBranchProbabilityInfoPassPass(PassRegistry&); void initializeLazyMachineBlockFrequencyInfoPassPass(PassRegistry&); void initializeLazyValueInfoPrinterPass(PassRegistry&); void initializeLazyValueInfoWrapperPassPass(PassRegistry&); +void initializeLegacyDivergenceAnalysisPass(PassRegistry&); void initializeLegacyLICMPassPass(PassRegistry&); void initializeLegacyLoopSinkPassPass(PassRegistry&); void initializeLegalizerPass(PassRegistry&); diff --git a/llvm/include/llvm/LinkAllPasses.h b/llvm/include/llvm/LinkAllPasses.h index bd432c5..f1de82d 100644 --- a/llvm/include/llvm/LinkAllPasses.h +++ b/llvm/include/llvm/LinkAllPasses.h @@ -94,7 +94,6 @@ namespace { (void) llvm::createDeadInstEliminationPass(); (void) llvm::createDeadStoreEliminationPass(); (void) llvm::createDependenceAnalysisWrapperPass(); - (void) llvm::createDivergenceAnalysisPass(); (void) llvm::createDomOnlyPrinterPass(); (void) llvm::createDomPrinterPass(); (void) llvm::createDomOnlyViewerPass(); @@ -121,6 +120,7 @@ namespace { (void) llvm::createInstructionCombiningPass(); (void) llvm::createInternalizePass(); (void) llvm::createLCSSAPass(); + (void) llvm::createLegacyDivergenceAnalysisPass(); (void) llvm::createLICMPass(); (void) llvm::createLoopSinkPass(); (void) llvm::createLazyValueInfoPass(); diff --git a/llvm/lib/Analysis/Analysis.cpp b/llvm/lib/Analysis/Analysis.cpp index 30576cf..49b1bdb 100644 --- a/llvm/lib/Analysis/Analysis.cpp +++ b/llvm/lib/Analysis/Analysis.cpp @@ -39,7 +39,6 @@ void llvm::initializeAnalysis(PassRegistry &Registry) { initializeDependenceAnalysisWrapperPassPass(Registry); initializeDelinearizationPass(Registry); initializeDemandedBitsWrapperPassPass(Registry); - initializeDivergenceAnalysisPass(Registry); initializeDominanceFrontierWrapperPassPass(Registry); initializeDomViewerPass(Registry); initializeDomPrinterPass(Registry); @@ -58,6 +57,7 @@ void llvm::initializeAnalysis(PassRegistry &Registry) { initializeLazyBlockFrequencyInfoPassPass(Registry); initializeLazyValueInfoWrapperPassPass(Registry); initializeLazyValueInfoPrinterPass(Registry); + initializeLegacyDivergenceAnalysisPass(Registry); initializeLintPass(Registry); initializeLoopInfoWrapperPassPass(Registry); initializeMemDepPrinterPass(Registry); diff --git a/llvm/lib/Analysis/CMakeLists.txt b/llvm/lib/Analysis/CMakeLists.txt index 6a228e8..5f20b20 100644 --- a/llvm/lib/Analysis/CMakeLists.txt +++ b/llvm/lib/Analysis/CMakeLists.txt @@ -25,7 +25,6 @@ add_llvm_library(LLVMAnalysis Delinearization.cpp DemandedBits.cpp DependenceAnalysis.cpp - DivergenceAnalysis.cpp DomPrinter.cpp DominanceFrontier.cpp EHPersonalities.cpp @@ -44,6 +43,7 @@ add_llvm_library(LLVMAnalysis LazyBlockFrequencyInfo.cpp LazyCallGraph.cpp LazyValueInfo.cpp + LegacyDivergenceAnalysis.cpp Lint.cpp Loads.cpp LoopAccessAnalysis.cpp diff --git a/llvm/lib/Analysis/DivergenceAnalysis.cpp b/llvm/lib/Analysis/LegacyDivergenceAnalysis.cpp similarity index 94% rename from llvm/lib/Analysis/DivergenceAnalysis.cpp rename to llvm/lib/Analysis/LegacyDivergenceAnalysis.cpp index f5f1874..c417862 100644 --- a/llvm/lib/Analysis/DivergenceAnalysis.cpp +++ b/llvm/lib/Analysis/LegacyDivergenceAnalysis.cpp @@ -1,4 +1,4 @@ -//===- DivergenceAnalysis.cpp --------- Divergence Analysis Implementation -==// +//===- LegacyDivergenceAnalysis.cpp --------- Legacy Divergence Analysis Implementation -==// // // The LLVM Compiler Infrastructure // @@ -64,7 +64,7 @@ // //===----------------------------------------------------------------------===// -#include "llvm/Analysis/DivergenceAnalysis.h" +#include "llvm/Analysis/LegacyDivergenceAnalysis.h" #include "llvm/Analysis/Passes.h" #include "llvm/Analysis/PostDominators.h" #include "llvm/Analysis/TargetTransformInfo.h" @@ -265,25 +265,25 @@ void DivergencePropagator::propagate() { } /// end namespace anonymous // Register this pass. -char DivergenceAnalysis::ID = 0; -INITIALIZE_PASS_BEGIN(DivergenceAnalysis, "divergence", "Divergence Analysis", +char LegacyDivergenceAnalysis::ID = 0; +INITIALIZE_PASS_BEGIN(LegacyDivergenceAnalysis, "divergence", "Legacy Divergence Analysis", false, true) INITIALIZE_PASS_DEPENDENCY(DominatorTreeWrapperPass) INITIALIZE_PASS_DEPENDENCY(PostDominatorTreeWrapperPass) -INITIALIZE_PASS_END(DivergenceAnalysis, "divergence", "Divergence Analysis", +INITIALIZE_PASS_END(LegacyDivergenceAnalysis, "divergence", "Legacy Divergence Analysis", false, true) -FunctionPass *llvm::createDivergenceAnalysisPass() { - return new DivergenceAnalysis(); +FunctionPass *llvm::createLegacyDivergenceAnalysisPass() { + return new LegacyDivergenceAnalysis(); } -void DivergenceAnalysis::getAnalysisUsage(AnalysisUsage &AU) const { +void LegacyDivergenceAnalysis::getAnalysisUsage(AnalysisUsage &AU) const { AU.addRequired(); AU.addRequired(); AU.setPreservesAll(); } -bool DivergenceAnalysis::runOnFunction(Function &F) { +bool LegacyDivergenceAnalysis::runOnFunction(Function &F) { auto *TTIWP = getAnalysisIfAvailable(); if (TTIWP == nullptr) return false; @@ -308,7 +308,7 @@ bool DivergenceAnalysis::runOnFunction(Function &F) { return false; } -void DivergenceAnalysis::print(raw_ostream &OS, const Module *) const { +void LegacyDivergenceAnalysis::print(raw_ostream &OS, const Module *) const { if (DivergentValues.empty()) return; const Value *FirstDivergentValue = *DivergentValues.begin(); diff --git a/llvm/lib/CodeGen/SelectionDAG/SelectionDAG.cpp b/llvm/lib/CodeGen/SelectionDAG/SelectionDAG.cpp index 944ec4c..19d569a 100644 --- a/llvm/lib/CodeGen/SelectionDAG/SelectionDAG.cpp +++ b/llvm/lib/CodeGen/SelectionDAG/SelectionDAG.cpp @@ -984,7 +984,7 @@ SelectionDAG::SelectionDAG(const TargetMachine &tm, CodeGenOpt::Level OL) void SelectionDAG::init(MachineFunction &NewMF, OptimizationRemarkEmitter &NewORE, Pass *PassPtr, const TargetLibraryInfo *LibraryInfo, - DivergenceAnalysis * Divergence) { + LegacyDivergenceAnalysis * Divergence) { MF = &NewMF; SDAGISelPass = PassPtr; ORE = &NewORE; diff --git a/llvm/lib/CodeGen/SelectionDAG/SelectionDAGISel.cpp b/llvm/lib/CodeGen/SelectionDAG/SelectionDAGISel.cpp index db4edb1..e104fae 100644 --- a/llvm/lib/CodeGen/SelectionDAG/SelectionDAGISel.cpp +++ b/llvm/lib/CodeGen/SelectionDAG/SelectionDAGISel.cpp @@ -417,7 +417,7 @@ bool SelectionDAGISel::runOnMachineFunction(MachineFunction &mf) { SplitCriticalSideEffectEdges(const_cast(Fn), DT, LI); CurDAG->init(*MF, *ORE, this, LibInfo, - getAnalysisIfAvailable()); + getAnalysisIfAvailable()); FuncInfo->set(Fn, *MF, CurDAG); // Now get the optional analyzes if we want to. diff --git a/llvm/lib/Target/AMDGPU/AMDGPUAnnotateUniformValues.cpp b/llvm/lib/Target/AMDGPU/AMDGPUAnnotateUniformValues.cpp index ed53708..d372271 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUAnnotateUniformValues.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUAnnotateUniformValues.cpp @@ -16,7 +16,7 @@ #include "AMDGPU.h" #include "AMDGPUIntrinsicInfo.h" #include "llvm/ADT/SetVector.h" -#include "llvm/Analysis/DivergenceAnalysis.h" +#include "llvm/Analysis/LegacyDivergenceAnalysis.h" #include "llvm/Analysis/LoopInfo.h" #include "llvm/Analysis/MemoryDependenceAnalysis.h" #include "llvm/IR/IRBuilder.h" @@ -32,7 +32,7 @@ namespace { class AMDGPUAnnotateUniformValues : public FunctionPass, public InstVisitor { - DivergenceAnalysis *DA; + LegacyDivergenceAnalysis *DA; MemoryDependenceResults *MDR; LoopInfo *LI; DenseMap noClobberClones; @@ -49,7 +49,7 @@ public: return "AMDGPU Annotate Uniform Values"; } void getAnalysisUsage(AnalysisUsage &AU) const override { - AU.addRequired(); + AU.addRequired(); AU.addRequired(); AU.addRequired(); AU.setPreservesAll(); @@ -64,7 +64,7 @@ public: INITIALIZE_PASS_BEGIN(AMDGPUAnnotateUniformValues, DEBUG_TYPE, "Add AMDGPU uniform metadata", false, false) -INITIALIZE_PASS_DEPENDENCY(DivergenceAnalysis) +INITIALIZE_PASS_DEPENDENCY(LegacyDivergenceAnalysis) INITIALIZE_PASS_DEPENDENCY(MemoryDependenceWrapperPass) INITIALIZE_PASS_DEPENDENCY(LoopInfoWrapperPass) INITIALIZE_PASS_END(AMDGPUAnnotateUniformValues, DEBUG_TYPE, @@ -176,7 +176,7 @@ bool AMDGPUAnnotateUniformValues::runOnFunction(Function &F) { if (skipFunction(F)) return false; - DA = &getAnalysis(); + DA = &getAnalysis(); MDR = &getAnalysis().getMemDep(); LI = &getAnalysis().getLoopInfo(); isKernelFunc = F.getCallingConv() == CallingConv::AMDGPU_KERNEL; diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCodeGenPrepare.cpp b/llvm/lib/Target/AMDGPU/AMDGPUCodeGenPrepare.cpp index 5713b7b..b821247 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUCodeGenPrepare.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUCodeGenPrepare.cpp @@ -18,7 +18,7 @@ #include "AMDGPUTargetMachine.h" #include "llvm/ADT/StringRef.h" #include "llvm/Analysis/AssumptionCache.h" -#include "llvm/Analysis/DivergenceAnalysis.h" +#include "llvm/Analysis/LegacyDivergenceAnalysis.h" #include "llvm/Analysis/Loads.h" #include "llvm/Analysis/ValueTracking.h" #include "llvm/CodeGen/Passes.h" @@ -60,7 +60,7 @@ class AMDGPUCodeGenPrepare : public FunctionPass, public InstVisitor { const GCNSubtarget *ST = nullptr; AssumptionCache *AC = nullptr; - DivergenceAnalysis *DA = nullptr; + LegacyDivergenceAnalysis *DA = nullptr; Module *Mod = nullptr; bool HasUnsafeFPMath = false; AMDGPUAS AMDGPUASI; @@ -177,7 +177,7 @@ public: void getAnalysisUsage(AnalysisUsage &AU) const override { AU.addRequired(); - AU.addRequired(); + AU.addRequired(); AU.setPreservesAll(); } }; @@ -898,7 +898,7 @@ bool AMDGPUCodeGenPrepare::runOnFunction(Function &F) { const AMDGPUTargetMachine &TM = TPC->getTM(); ST = &TM.getSubtarget(F); AC = &getAnalysis().getAssumptionCache(F); - DA = &getAnalysis(); + DA = &getAnalysis(); HasUnsafeFPMath = hasUnsafeFPMath(F); AMDGPUASI = TM.getAMDGPUAS(); @@ -918,7 +918,7 @@ bool AMDGPUCodeGenPrepare::runOnFunction(Function &F) { INITIALIZE_PASS_BEGIN(AMDGPUCodeGenPrepare, DEBUG_TYPE, "AMDGPU IR optimizations", false, false) INITIALIZE_PASS_DEPENDENCY(AssumptionCacheTracker) -INITIALIZE_PASS_DEPENDENCY(DivergenceAnalysis) +INITIALIZE_PASS_DEPENDENCY(LegacyDivergenceAnalysis) INITIALIZE_PASS_END(AMDGPUCodeGenPrepare, DEBUG_TYPE, "AMDGPU IR optimizations", false, false) diff --git a/llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp b/llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp index e70de24..f4af9cd 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp @@ -29,7 +29,7 @@ #include "llvm/ADT/APInt.h" #include "llvm/ADT/SmallVector.h" #include "llvm/ADT/StringRef.h" -#include "llvm/Analysis/DivergenceAnalysis.h" +#include "llvm/Analysis/LegacyDivergenceAnalysis.h" #include "llvm/Analysis/ValueTracking.h" #include "llvm/CodeGen/FunctionLoweringInfo.h" #include "llvm/CodeGen/ISDOpcodes.h" @@ -87,7 +87,7 @@ public: void getAnalysisUsage(AnalysisUsage &AU) const override { AU.addRequired(); AU.addRequired(); - AU.addRequired(); + AU.addRequired(); SelectionDAGISel::getAnalysisUsage(AU); } @@ -253,7 +253,7 @@ INITIALIZE_PASS_BEGIN(AMDGPUDAGToDAGISel, "isel", "AMDGPU DAG->DAG Pattern Instruction Selection", false, false) INITIALIZE_PASS_DEPENDENCY(AMDGPUArgumentUsageInfo) INITIALIZE_PASS_DEPENDENCY(AMDGPUPerfHintAnalysis) -INITIALIZE_PASS_DEPENDENCY(DivergenceAnalysis) +INITIALIZE_PASS_DEPENDENCY(LegacyDivergenceAnalysis) INITIALIZE_PASS_END(AMDGPUDAGToDAGISel, "isel", "AMDGPU DAG->DAG Pattern Instruction Selection", false, false) diff --git a/llvm/lib/Target/AMDGPU/AMDGPULowerKernelArguments.cpp b/llvm/lib/Target/AMDGPU/AMDGPULowerKernelArguments.cpp index c147830..2a25f1a 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPULowerKernelArguments.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPULowerKernelArguments.cpp @@ -16,7 +16,6 @@ #include "AMDGPUSubtarget.h" #include "AMDGPUTargetMachine.h" #include "llvm/ADT/StringRef.h" -#include "llvm/Analysis/DivergenceAnalysis.h" #include "llvm/Analysis/Loads.h" #include "llvm/CodeGen/Passes.h" #include "llvm/CodeGen/TargetPassConfig.h" diff --git a/llvm/lib/Target/AMDGPU/AMDGPUUnifyDivergentExitNodes.cpp b/llvm/lib/Target/AMDGPU/AMDGPUUnifyDivergentExitNodes.cpp index 0d3a167..ced3f6f 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUUnifyDivergentExitNodes.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUUnifyDivergentExitNodes.cpp @@ -25,7 +25,7 @@ #include "llvm/ADT/SmallPtrSet.h" #include "llvm/ADT/SmallVector.h" #include "llvm/ADT/StringRef.h" -#include "llvm/Analysis/DivergenceAnalysis.h" +#include "llvm/Analysis/LegacyDivergenceAnalysis.h" #include "llvm/Analysis/PostDominators.h" #include "llvm/Analysis/TargetTransformInfo.h" #include "llvm/Transforms/Utils/Local.h" @@ -70,7 +70,7 @@ char &llvm::AMDGPUUnifyDivergentExitNodesID = AMDGPUUnifyDivergentExitNodes::ID; INITIALIZE_PASS_BEGIN(AMDGPUUnifyDivergentExitNodes, DEBUG_TYPE, "Unify divergent function exit nodes", false, false) INITIALIZE_PASS_DEPENDENCY(PostDominatorTreeWrapperPass) -INITIALIZE_PASS_DEPENDENCY(DivergenceAnalysis) +INITIALIZE_PASS_DEPENDENCY(LegacyDivergenceAnalysis) INITIALIZE_PASS_END(AMDGPUUnifyDivergentExitNodes, DEBUG_TYPE, "Unify divergent function exit nodes", false, false) @@ -78,10 +78,10 @@ void AMDGPUUnifyDivergentExitNodes::getAnalysisUsage(AnalysisUsage &AU) const{ // TODO: Preserve dominator tree. AU.addRequired(); - AU.addRequired(); + AU.addRequired(); // No divergent values are changed, only blocks and branch edges. - AU.addPreserved(); + AU.addPreserved(); // We preserve the non-critical-edgeness property AU.addPreservedID(BreakCriticalEdgesID); @@ -95,7 +95,7 @@ void AMDGPUUnifyDivergentExitNodes::getAnalysisUsage(AnalysisUsage &AU) const{ /// \returns true if \p BB is reachable through only uniform branches. /// XXX - Is there a more efficient way to find this? -static bool isUniformlyReached(const DivergenceAnalysis &DA, +static bool isUniformlyReached(const LegacyDivergenceAnalysis &DA, BasicBlock &BB) { SmallVector Stack; SmallPtrSet Visited; @@ -163,7 +163,7 @@ bool AMDGPUUnifyDivergentExitNodes::runOnFunction(Function &F) { if (PDT.getRoots().size() <= 1) return false; - DivergenceAnalysis &DA = getAnalysis(); + LegacyDivergenceAnalysis &DA = getAnalysis(); // Loop over all of the blocks in a function, tracking all of the blocks that // return. diff --git a/llvm/lib/Target/AMDGPU/SIAnnotateControlFlow.cpp b/llvm/lib/Target/AMDGPU/SIAnnotateControlFlow.cpp index a23e102..c52313f 100644 --- a/llvm/lib/Target/AMDGPU/SIAnnotateControlFlow.cpp +++ b/llvm/lib/Target/AMDGPU/SIAnnotateControlFlow.cpp @@ -16,7 +16,7 @@ #include "llvm/ADT/DepthFirstIterator.h" #include "llvm/ADT/STLExtras.h" #include "llvm/ADT/SmallVector.h" -#include "llvm/Analysis/DivergenceAnalysis.h" +#include "llvm/Analysis/LegacyDivergenceAnalysis.h" #include "llvm/Analysis/LoopInfo.h" #include "llvm/Transforms/Utils/Local.h" #include "llvm/IR/BasicBlock.h" @@ -52,7 +52,7 @@ using StackEntry = std::pair; using StackVector = SmallVector; class SIAnnotateControlFlow : public FunctionPass { - DivergenceAnalysis *DA; + LegacyDivergenceAnalysis *DA; Type *Boolean; Type *Void; @@ -116,7 +116,7 @@ public: void getAnalysisUsage(AnalysisUsage &AU) const override { AU.addRequired(); AU.addRequired(); - AU.addRequired(); + AU.addRequired(); AU.addPreserved(); FunctionPass::getAnalysisUsage(AU); } @@ -127,7 +127,7 @@ public: INITIALIZE_PASS_BEGIN(SIAnnotateControlFlow, DEBUG_TYPE, "Annotate SI Control Flow", false, false) INITIALIZE_PASS_DEPENDENCY(DominatorTreeWrapperPass) -INITIALIZE_PASS_DEPENDENCY(DivergenceAnalysis) +INITIALIZE_PASS_DEPENDENCY(LegacyDivergenceAnalysis) INITIALIZE_PASS_END(SIAnnotateControlFlow, DEBUG_TYPE, "Annotate SI Control Flow", false, false) @@ -387,7 +387,7 @@ void SIAnnotateControlFlow::closeControlFlow(BasicBlock *BB) { bool SIAnnotateControlFlow::runOnFunction(Function &F) { DT = &getAnalysis().getDomTree(); LI = &getAnalysis().getLoopInfo(); - DA = &getAnalysis(); + DA = &getAnalysis(); for (df_iterator I = df_begin(&F.getEntryBlock()), E = df_end(&F.getEntryBlock()); I != E; ++I) { diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp index 5916395..9f4599f 100644 --- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp +++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp @@ -9166,7 +9166,7 @@ void SITargetLowering::computeKnownBitsForFrameIndex(const SDValue Op, } bool SITargetLowering::isSDNodeSourceOfDivergence(const SDNode * N, - FunctionLoweringInfo * FLI, DivergenceAnalysis * DA) const + FunctionLoweringInfo * FLI, LegacyDivergenceAnalysis * KDA) const { switch (N->getOpcode()) { case ISD::Register: @@ -9199,7 +9199,7 @@ bool SITargetLowering::isSDNodeSourceOfDivergence(const SDNode * N, else if (!AMDGPU::isEntryFunctionCC(FLI->Fn->getCallingConv())) return true; } - return !DA || DA->isDivergent(FLI->getValueFromVirtualReg(Reg)); + return !KDA || KDA->isDivergent(FLI->getValueFromVirtualReg(Reg)); } } break; diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.h b/llvm/lib/Target/AMDGPU/SIISelLowering.h index b426537..6c02688 100644 --- a/llvm/lib/Target/AMDGPU/SIISelLowering.h +++ b/llvm/lib/Target/AMDGPU/SIISelLowering.h @@ -339,7 +339,7 @@ public: unsigned Depth = 0) const override; bool isSDNodeSourceOfDivergence(const SDNode *N, - FunctionLoweringInfo *FLI, DivergenceAnalysis *DA) const override; + FunctionLoweringInfo *FLI, LegacyDivergenceAnalysis *DA) const override; bool isCanonicalized(SelectionDAG &DAG, SDValue Op, unsigned MaxDepth = 5) const; diff --git a/llvm/lib/Transforms/Scalar/LoopUnswitch.cpp b/llvm/lib/Transforms/Scalar/LoopUnswitch.cpp index 7169aed..7531cc0 100644 --- a/llvm/lib/Transforms/Scalar/LoopUnswitch.cpp +++ b/llvm/lib/Transforms/Scalar/LoopUnswitch.cpp @@ -33,7 +33,7 @@ #include "llvm/ADT/Statistic.h" #include "llvm/Analysis/AssumptionCache.h" #include "llvm/Analysis/CodeMetrics.h" -#include "llvm/Analysis/DivergenceAnalysis.h" +#include "llvm/Analysis/LegacyDivergenceAnalysis.h" #include "llvm/Analysis/InstructionSimplify.h" #include "llvm/Analysis/LoopInfo.h" #include "llvm/Analysis/LoopPass.h" @@ -215,7 +215,7 @@ namespace { AU.addRequired(); AU.addRequired(); if (hasBranchDivergence) - AU.addRequired(); + AU.addRequired(); getLoopAnalysisUsage(AU); } @@ -383,7 +383,7 @@ INITIALIZE_PASS_BEGIN(LoopUnswitch, "loop-unswitch", "Unswitch loops", INITIALIZE_PASS_DEPENDENCY(AssumptionCacheTracker) INITIALIZE_PASS_DEPENDENCY(LoopPass) INITIALIZE_PASS_DEPENDENCY(TargetTransformInfoWrapperPass) -INITIALIZE_PASS_DEPENDENCY(DivergenceAnalysis) +INITIALIZE_PASS_DEPENDENCY(LegacyDivergenceAnalysis) INITIALIZE_PASS_END(LoopUnswitch, "loop-unswitch", "Unswitch loops", false, false) @@ -864,7 +864,7 @@ bool LoopUnswitch::UnswitchIfProfitable(Value *LoopCond, Constant *Val, return false; } if (hasBranchDivergence && - getAnalysis().isDivergent(LoopCond)) { + getAnalysis().isDivergent(LoopCond)) { LLVM_DEBUG(dbgs() << "NOT unswitching loop %" << currentLoop->getHeader()->getName() << " at non-trivial condition '" << *Val diff --git a/llvm/lib/Transforms/Scalar/StructurizeCFG.cpp b/llvm/lib/Transforms/Scalar/StructurizeCFG.cpp index d650264..f58f79f 100644 --- a/llvm/lib/Transforms/Scalar/StructurizeCFG.cpp +++ b/llvm/lib/Transforms/Scalar/StructurizeCFG.cpp @@ -13,7 +13,7 @@ #include "llvm/ADT/STLExtras.h" #include "llvm/ADT/SmallPtrSet.h" #include "llvm/ADT/SmallVector.h" -#include "llvm/Analysis/DivergenceAnalysis.h" +#include "llvm/Analysis/LegacyDivergenceAnalysis.h" #include "llvm/Analysis/LoopInfo.h" #include "llvm/Analysis/RegionInfo.h" #include "llvm/Analysis/RegionIterator.h" @@ -183,7 +183,7 @@ class StructurizeCFG : public RegionPass { Function *Func; Region *ParentRegion; - DivergenceAnalysis *DA; + LegacyDivergenceAnalysis *DA; DominatorTree *DT; LoopInfo *LI; @@ -269,7 +269,7 @@ public: void getAnalysisUsage(AnalysisUsage &AU) const override { if (SkipUniformRegions) - AU.addRequired(); + AU.addRequired(); AU.addRequiredID(LowerSwitchID); AU.addRequired(); AU.addRequired(); @@ -285,7 +285,7 @@ char StructurizeCFG::ID = 0; INITIALIZE_PASS_BEGIN(StructurizeCFG, "structurizecfg", "Structurize the CFG", false, false) -INITIALIZE_PASS_DEPENDENCY(DivergenceAnalysis) +INITIALIZE_PASS_DEPENDENCY(LegacyDivergenceAnalysis) INITIALIZE_PASS_DEPENDENCY(LowerSwitch) INITIALIZE_PASS_DEPENDENCY(DominatorTreeWrapperPass) INITIALIZE_PASS_DEPENDENCY(RegionInfoPass) @@ -914,7 +914,7 @@ void StructurizeCFG::rebuildSSA() { } static bool hasOnlyUniformBranches(Region *R, unsigned UniformMDKindID, - const DivergenceAnalysis &DA) { + const LegacyDivergenceAnalysis &DA) { for (auto E : R->elements()) { if (!E->isSubRegion()) { auto Br = dyn_cast(E->getEntry()->getTerminator()); @@ -962,7 +962,7 @@ bool StructurizeCFG::runOnRegion(Region *R, RGPassManager &RGM) { // but we shouldn't rely on metadata for correctness! unsigned UniformMDKindID = R->getEntry()->getContext().getMDKindID("structurizecfg.uniform"); - DA = &getAnalysis(); + DA = &getAnalysis(); if (hasOnlyUniformBranches(R, UniformMDKindID, *DA)) { LLVM_DEBUG(dbgs() << "Skipping region with uniform control flow: " << *R diff --git a/llvm/test/Analysis/DivergenceAnalysis/AMDGPU/atomics.ll b/llvm/test/Analysis/LegacyDivergenceAnalysis/AMDGPU/atomics.ll similarity index 100% rename from llvm/test/Analysis/DivergenceAnalysis/AMDGPU/atomics.ll rename to llvm/test/Analysis/LegacyDivergenceAnalysis/AMDGPU/atomics.ll diff --git a/llvm/test/Analysis/DivergenceAnalysis/AMDGPU/intrinsics.ll b/llvm/test/Analysis/LegacyDivergenceAnalysis/AMDGPU/intrinsics.ll similarity index 100% rename from llvm/test/Analysis/DivergenceAnalysis/AMDGPU/intrinsics.ll rename to llvm/test/Analysis/LegacyDivergenceAnalysis/AMDGPU/intrinsics.ll diff --git a/llvm/test/Analysis/DivergenceAnalysis/AMDGPU/kernel-args.ll b/llvm/test/Analysis/LegacyDivergenceAnalysis/AMDGPU/kernel-args.ll similarity index 80% rename from llvm/test/Analysis/DivergenceAnalysis/AMDGPU/kernel-args.ll rename to llvm/test/Analysis/LegacyDivergenceAnalysis/AMDGPU/kernel-args.ll index 8ae4438..22f28ac 100644 --- a/llvm/test/Analysis/DivergenceAnalysis/AMDGPU/kernel-args.ll +++ b/llvm/test/Analysis/LegacyDivergenceAnalysis/AMDGPU/kernel-args.ll @@ -1,6 +1,6 @@ ; RUN: opt %s -mtriple amdgcn-- -analyze -divergence | FileCheck %s -; CHECK-LABEL: Printing analysis 'Divergence Analysis' for function 'test_amdgpu_ps': +; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'test_amdgpu_ps': ; CHECK: DIVERGENT: ; CHECK-NOT: %arg0 ; CHECK-NOT: %arg1 @@ -14,7 +14,7 @@ define amdgpu_ps void @test_amdgpu_ps([4 x <16 x i8>] addrspace(2)* byval %arg0, ret void } -; CHECK-LABEL: Printing analysis 'Divergence Analysis' for function 'test_amdgpu_kernel': +; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'test_amdgpu_kernel': ; CHECK-NOT: %arg0 ; CHECK-NOT: %arg1 ; CHECK-NOT: %arg2 @@ -26,7 +26,7 @@ define amdgpu_kernel void @test_amdgpu_kernel([4 x <16 x i8>] addrspace(2)* byva ret void } -; CHECK-LABEL: Printing analysis 'Divergence Analysis' for function 'test_c': +; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'test_c': ; CHECK: DIVERGENT: ; CHECK: DIVERGENT: ; CHECK: DIVERGENT: diff --git a/llvm/test/Analysis/DivergenceAnalysis/AMDGPU/lit.local.cfg b/llvm/test/Analysis/LegacyDivergenceAnalysis/AMDGPU/lit.local.cfg similarity index 100% rename from llvm/test/Analysis/DivergenceAnalysis/AMDGPU/lit.local.cfg rename to llvm/test/Analysis/LegacyDivergenceAnalysis/AMDGPU/lit.local.cfg diff --git a/llvm/test/Analysis/DivergenceAnalysis/AMDGPU/llvm.amdgcn.buffer.atomic.ll b/llvm/test/Analysis/LegacyDivergenceAnalysis/AMDGPU/llvm.amdgcn.buffer.atomic.ll similarity index 100% rename from llvm/test/Analysis/DivergenceAnalysis/AMDGPU/llvm.amdgcn.buffer.atomic.ll rename to llvm/test/Analysis/LegacyDivergenceAnalysis/AMDGPU/llvm.amdgcn.buffer.atomic.ll diff --git a/llvm/test/Analysis/DivergenceAnalysis/AMDGPU/llvm.amdgcn.image.atomic.ll b/llvm/test/Analysis/LegacyDivergenceAnalysis/AMDGPU/llvm.amdgcn.image.atomic.ll similarity index 100% rename from llvm/test/Analysis/DivergenceAnalysis/AMDGPU/llvm.amdgcn.image.atomic.ll rename to llvm/test/Analysis/LegacyDivergenceAnalysis/AMDGPU/llvm.amdgcn.image.atomic.ll diff --git a/llvm/test/Analysis/DivergenceAnalysis/AMDGPU/no-return-blocks.ll b/llvm/test/Analysis/LegacyDivergenceAnalysis/AMDGPU/no-return-blocks.ll similarity index 100% rename from llvm/test/Analysis/DivergenceAnalysis/AMDGPU/no-return-blocks.ll rename to llvm/test/Analysis/LegacyDivergenceAnalysis/AMDGPU/no-return-blocks.ll diff --git a/llvm/test/Analysis/DivergenceAnalysis/AMDGPU/phi-undef.ll b/llvm/test/Analysis/LegacyDivergenceAnalysis/AMDGPU/phi-undef.ll similarity index 100% rename from llvm/test/Analysis/DivergenceAnalysis/AMDGPU/phi-undef.ll rename to llvm/test/Analysis/LegacyDivergenceAnalysis/AMDGPU/phi-undef.ll diff --git a/llvm/test/Analysis/DivergenceAnalysis/AMDGPU/unreachable-loop-block.ll b/llvm/test/Analysis/LegacyDivergenceAnalysis/AMDGPU/unreachable-loop-block.ll similarity index 100% rename from llvm/test/Analysis/DivergenceAnalysis/AMDGPU/unreachable-loop-block.ll rename to llvm/test/Analysis/LegacyDivergenceAnalysis/AMDGPU/unreachable-loop-block.ll diff --git a/llvm/test/Analysis/DivergenceAnalysis/AMDGPU/workitem-intrinsics.ll b/llvm/test/Analysis/LegacyDivergenceAnalysis/AMDGPU/workitem-intrinsics.ll similarity index 100% rename from llvm/test/Analysis/DivergenceAnalysis/AMDGPU/workitem-intrinsics.ll rename to llvm/test/Analysis/LegacyDivergenceAnalysis/AMDGPU/workitem-intrinsics.ll diff --git a/llvm/test/Analysis/DivergenceAnalysis/NVPTX/diverge.ll b/llvm/test/Analysis/LegacyDivergenceAnalysis/NVPTX/diverge.ll similarity index 90% rename from llvm/test/Analysis/DivergenceAnalysis/NVPTX/diverge.ll rename to llvm/test/Analysis/LegacyDivergenceAnalysis/NVPTX/diverge.ll index cbfd290..4e7163d 100644 --- a/llvm/test/Analysis/DivergenceAnalysis/NVPTX/diverge.ll +++ b/llvm/test/Analysis/LegacyDivergenceAnalysis/NVPTX/diverge.ll @@ -5,7 +5,7 @@ target triple = "nvptx64-nvidia-cuda" ; return (n < 0 ? a + threadIdx.x : b + threadIdx.x) define i32 @no_diverge(i32 %n, i32 %a, i32 %b) { -; CHECK-LABEL: Printing analysis 'Divergence Analysis' for function 'no_diverge' +; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'no_diverge' entry: %tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() %cond = icmp slt i32 %n, 0 @@ -27,7 +27,7 @@ merge: ; c = b; ; return c; // c is divergent: sync dependent define i32 @sync(i32 %a, i32 %b) { -; CHECK-LABEL: Printing analysis 'Divergence Analysis' for function 'sync' +; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'sync' bb1: %tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.y() %cond = icmp slt i32 %tid, 5 @@ -48,7 +48,7 @@ bb3: ; // c here is divergent because it is sync dependent on threadIdx.x >= 5 ; return c; define i32 @mixed(i32 %n, i32 %a, i32 %b) { -; CHECK-LABEL: Printing analysis 'Divergence Analysis' for function 'mixed' +; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'mixed' bb1: %tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.z() %cond = icmp slt i32 %tid, 5 @@ -73,7 +73,7 @@ bb6: ; We conservatively treats all parameters of a __device__ function as divergent. define i32 @device(i32 %n, i32 %a, i32 %b) { -; CHECK-LABEL: Printing analysis 'Divergence Analysis' for function 'device' +; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'device' ; CHECK: DIVERGENT: i32 %n ; CHECK: DIVERGENT: i32 %a ; CHECK: DIVERGENT: i32 %b @@ -98,7 +98,7 @@ merge: ; ; The i defined in the loop is used outside. define i32 @loop() { -; CHECK-LABEL: Printing analysis 'Divergence Analysis' for function 'loop' +; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'loop' entry: %laneid = call i32 @llvm.nvvm.read.ptx.sreg.laneid() br label %loop @@ -120,7 +120,7 @@ else: ; Same as @loop, but the loop is in the LCSSA form. define i32 @lcssa() { -; CHECK-LABEL: Printing analysis 'Divergence Analysis' for function 'lcssa' +; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'lcssa' entry: %tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() br label %loop @@ -156,7 +156,7 @@ else: ; if (i3 == 5) // divergent ; because sync dependent on (tid / i3). define i32 @unstructured_loop(i1 %entry_cond) { -; CHECK-LABEL: Printing analysis 'Divergence Analysis' for function 'unstructured_loop' +; CHECK-LABEL: Printing analysis 'Legacy Divergence Analysis' for function 'unstructured_loop' entry: %tid = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() br i1 %entry_cond, label %loop_entry_1, label %loop_entry_2 diff --git a/llvm/test/Analysis/DivergenceAnalysis/NVPTX/lit.local.cfg b/llvm/test/Analysis/LegacyDivergenceAnalysis/NVPTX/lit.local.cfg similarity index 100% rename from llvm/test/Analysis/DivergenceAnalysis/NVPTX/lit.local.cfg rename to llvm/test/Analysis/LegacyDivergenceAnalysis/NVPTX/lit.local.cfg -- 2.7.4