[OpenMP] Add OpenMPOpt as a Module pass

[jhuber6]

Summary

This patch registers OpenMPOpt as a Module pass in addition to a CGSCC
pass. This is so certain optimzations that are sensitive to intact
call-sites can happen before inlining. The old openmpopt pass name is
changed to openmp-opt-cgscc and openmp-opt calls the Module pass
which has limited functionality currently.

[lebedev.ri]

This needs PhaseOrdering test

[jhuber6]

Forgot to format.

[jhuber6]

Changed the pass registration to happen before the Attributor pass to avoid function specialization. Added a PhaseOrdering test that won't pass without the module pass.

[jdoerfert]

Can you precommit the llvm/test/Transforms/PhaseOrdering/openmp-opt-module.ll test so we can see what changed.

Generally looks good to me otherwise.

llvm/test/Transforms/OpenMP/values_in_offload_arrays.ll

1–2

make two run lines, here and elsewhere if applicable

[jhuber6]

Changing the OpenMPOpt module pass to only run selected optimizations. Removed the module pass version from most of the tests, more tests will be added once there are additional uses for the module pass.

[jhuber6]

Rebase.

[jhuber6]

Fixing tests.

[mnadeem]

The following two tests are causing assertion failures when run with the legacy pass manager (force using -enable-new-pm=0).

• Transforms/OpenMP/hide_mem_transfer_latency.ll
• Transforms/OpenMP/values_in_offload_arrays.ll

Waiting on the bugzilla account request to make a ticket but here is the trace.

opt: llvm/lib/Analysis/CallGraphSCCPass.cpp:281: bool (anonymous namespace)::CGPassManager::RefreshCallGraph(const llvm::CallGraphSCC &, llvm::CallGraph &, bool): Assertion `!CheckingMode && "CallGraphSCCPass did not update the CallGraph correctly!"' failed.
PLEASE submit a bug report to https://bugs.llvm.org/ and include the crash backtrace.
Stack dump:
0.      Program arguments: opt -S -openmp-opt-cgscc -aa-pipeline=basic-aa -openmp-hide-memory-transfer-latency -enable-new-pm=0
1.      Running pass 'CallGraph Pass Manager' on module '<stdin>'.
 #0 0x00000000020bdf43 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (opt+0x20bdf43)
 #1 0x00000000020bbc00 llvm::sys::RunSignalHandlers() (opt+0x20bbc00)
 #2 0x00000000020be3f6 SignalHandler(int) llvm/lib/Support/Signals.cpp:0:0
 #3 0x00007f15159bf330 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x10330)
 #4 0x00007f15143a3c37 raise /build/eglibc-ripdx6/eglibc-2.19/signal/../nptl/sysdeps/unix/sysv/linux/raise.c:56:0
 #5 0x00007f15143a7028 abort /build/eglibc-ripdx6/eglibc-2.19/stdlib/abort.c:91:0
 #6 0x00007f151439cbf6 __assert_fail_base /build/eglibc-ripdx6/eglibc-2.19/assert/assert.c:92:0
 #7 0x00007f151439cca2 (/lib/x86_64-linux-gnu/libc.so.6+0x2fca2)
 #8 0x000000000100ab4f (anonymous namespace)::CGPassManager::RefreshCallGraph(llvm::CallGraphSCC const&, llvm::CallGraph&, bool) llvm/lib/Analysis/CallGraphSCCPass.cpp:0:0
 #9 0x00000000010093ac (anonymous namespace)::CGPassManager::runOnModule(llvm::Module&) llvm/lib/Analysis/CallGraphSCCPass.cpp:0:0
#10 0x00000000018c6788 llvm::legacy::PassManagerImpl::run(llvm::Module&) (opt+0x18c6788)
#11 0x00000000007d06e5 main (opt+0x7d06e5)
#12 0x00007f151438ef45 __libc_start_main /build/eglibc-ripdx6/eglibc-2.19/csu/libc-start.c:321:0
#13 0x00000000007ba679 _start (opt+0x7ba679)
[1]    27229 abort      opt -S -openmp-opt-cgscc -aa-pipeline=basic-aa   -enable-new-pm=0

[jhuber6]

In D99202#2787617, @mnadeem wrote:

The following two tests are causing assertion failures when run with the legacy pass manager (force using -enable-new-pm=0).

• Transforms/OpenMP/hide_mem_transfer_latency.ll
• Transforms/OpenMP/values_in_offload_arrays.ll

Waiting on the bugzilla account request to make a ticket but here is the trace.

opt: llvm/lib/Analysis/CallGraphSCCPass.cpp:281: bool (anonymous namespace)::CGPassManager::RefreshCallGraph(const llvm::CallGraphSCC &, llvm::CallGraph &, bool): Assertion `!CheckingMode && "CallGraphSCCPass did not update the CallGraph correctly!"' failed.
PLEASE submit a bug report to https://bugs.llvm.org/ and include the crash backtrace.
Stack dump:
0.      Program arguments: opt -S -openmp-opt-cgscc -aa-pipeline=basic-aa -openmp-hide-memory-transfer-latency -enable-new-pm=0
1.      Running pass 'CallGraph Pass Manager' on module '<stdin>'.
 #0 0x00000000020bdf43 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (opt+0x20bdf43)
 #1 0x00000000020bbc00 llvm::sys::RunSignalHandlers() (opt+0x20bbc00)
 #2 0x00000000020be3f6 SignalHandler(int) llvm/lib/Support/Signals.cpp:0:0
 #3 0x00007f15159bf330 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x10330)
 #4 0x00007f15143a3c37 raise /build/eglibc-ripdx6/eglibc-2.19/signal/../nptl/sysdeps/unix/sysv/linux/raise.c:56:0
 #5 0x00007f15143a7028 abort /build/eglibc-ripdx6/eglibc-2.19/stdlib/abort.c:91:0
 #6 0x00007f151439cbf6 __assert_fail_base /build/eglibc-ripdx6/eglibc-2.19/assert/assert.c:92:0
 #7 0x00007f151439cca2 (/lib/x86_64-linux-gnu/libc.so.6+0x2fca2)
 #8 0x000000000100ab4f (anonymous namespace)::CGPassManager::RefreshCallGraph(llvm::CallGraphSCC const&, llvm::CallGraph&, bool) llvm/lib/Analysis/CallGraphSCCPass.cpp:0:0
 #9 0x00000000010093ac (anonymous namespace)::CGPassManager::runOnModule(llvm::Module&) llvm/lib/Analysis/CallGraphSCCPass.cpp:0:0
#10 0x00000000018c6788 llvm::legacy::PassManagerImpl::run(llvm::Module&) (opt+0x18c6788)
#11 0x00000000007d06e5 main (opt+0x7d06e5)
#12 0x00007f151438ef45 __libc_start_main /build/eglibc-ripdx6/eglibc-2.19/csu/libc-start.c:321:0
#13 0x00000000007ba679 _start (opt+0x7ba679)
[1]    27229 abort      opt -S -openmp-opt-cgscc -aa-pipeline=basic-aa   -enable-new-pm=0

I'm not sure what the cause of this is, I committed rGb19136e35256c5b21aff93ea5b55b4c142408c40 as a quick work-around. I don't think this patch is the cause, I checked the revision prior to this one landing and it had the same issue. I'm not sure how much of an issue it is overall considering that the new pass manager is the default and hiding the memory transfer latency is an optional feature.

[jdoerfert]

In D99202#2787617, @mnadeem wrote:

The following two tests are causing assertion failures when run with the legacy pass manager (force using -enable-new-pm=0).

• Transforms/OpenMP/hide_mem_transfer_latency.ll
• Transforms/OpenMP/values_in_offload_arrays.ll

Waiting on the bugzilla account request to make a ticket but here is the trace.

opt: llvm/lib/Analysis/CallGraphSCCPass.cpp:281: bool (anonymous namespace)::CGPassManager::RefreshCallGraph(const llvm::CallGraphSCC &, llvm::CallGraph &, bool): Assertion `!CheckingMode && "CallGraphSCCPass did not update the CallGraph correctly!"' failed.
PLEASE submit a bug report to https://bugs.llvm.org/ and include the crash backtrace.
Stack dump:
0.      Program arguments: opt -S -openmp-opt-cgscc -aa-pipeline=basic-aa -openmp-hide-memory-transfer-latency -enable-new-pm=0
1.      Running pass 'CallGraph Pass Manager' on module '<stdin>'.
 #0 0x00000000020bdf43 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (opt+0x20bdf43)
 #1 0x00000000020bbc00 llvm::sys::RunSignalHandlers() (opt+0x20bbc00)
 #2 0x00000000020be3f6 SignalHandler(int) llvm/lib/Support/Signals.cpp:0:0
 #3 0x00007f15159bf330 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x10330)
 #4 0x00007f15143a3c37 raise /build/eglibc-ripdx6/eglibc-2.19/signal/../nptl/sysdeps/unix/sysv/linux/raise.c:56:0
 #5 0x00007f15143a7028 abort /build/eglibc-ripdx6/eglibc-2.19/stdlib/abort.c:91:0
 #6 0x00007f151439cbf6 __assert_fail_base /build/eglibc-ripdx6/eglibc-2.19/assert/assert.c:92:0
 #7 0x00007f151439cca2 (/lib/x86_64-linux-gnu/libc.so.6+0x2fca2)
 #8 0x000000000100ab4f (anonymous namespace)::CGPassManager::RefreshCallGraph(llvm::CallGraphSCC const&, llvm::CallGraph&, bool) llvm/lib/Analysis/CallGraphSCCPass.cpp:0:0
 #9 0x00000000010093ac (anonymous namespace)::CGPassManager::runOnModule(llvm::Module&) llvm/lib/Analysis/CallGraphSCCPass.cpp:0:0
#10 0x00000000018c6788 llvm::legacy::PassManagerImpl::run(llvm::Module&) (opt+0x18c6788)
#11 0x00000000007d06e5 main (opt+0x7d06e5)
#12 0x00007f151438ef45 __libc_start_main /build/eglibc-ripdx6/eglibc-2.19/csu/libc-start.c:321:0
#13 0x00000000007ba679 _start (opt+0x7ba679)
[1]    27229 abort      opt -S -openmp-opt-cgscc -aa-pipeline=basic-aa   -enable-new-pm=0

Attributor bugs, I thought I fixed them. Do they still happen in ToT?

llvm/include/llvm/InitializePasses.h

 void initializeAggressiveInstCombinerLegacyPassPass(PassRegistry&);
 void initializeAliasSetPrinterPass(PassRegistry&);
 void initializeAlignmentFromAssumptionsPass(PassRegistry&);
 void initializeAlwaysInlinerLegacyPassPass(PassRegistry&);
 void initializeAssumeSimplifyPassLegacyPassPass(PassRegistry &);
 void initializeAssumeBuilderPassLegacyPassPass(PassRegistry &);
 void initializeAnnotation2MetadataLegacyPass(PassRegistry &);
 void initializeAnnotationRemarksLegacyPass(PassRegistry &);
-void initializeOpenMPOptLegacyPassPass(PassRegistry &);
+void initializeOpenMPOptCGSCCLegacyPassPass(PassRegistry &);
 void initializeArgPromotionPass(PassRegistry&);
 void initializeAssumptionCacheTrackerPass(PassRegistry&);
 void initializeAtomicExpandPass(PassRegistry&);
 void initializeAttributorLegacyPassPass(PassRegistry&);
 void initializeAttributorCGSCCLegacyPassPass(PassRegistry &);
 void initializeBasicBlockSectionsPass(PassRegistry &);
 void initializeBDCELegacyPassPass(PassRegistry&);
 void initializeBarrierNoopPass(PassRegistry&);

llvm/include/llvm/LinkAllPasses.h

ForcePassLinking() {
       // to know that getenv() never returns -1, this will do the job.
       if (std::getenv("bar") != (char*) -1)
         return;
 
       (void) llvm::createAAEvalPass();
       (void) llvm::createAggressiveDCEPass();
       (void) llvm::createAggressiveInstCombinerPass();
       (void) llvm::createBitTrackingDCEPass();
-      (void) llvm::createOpenMPOptLegacyPass();
+      (void) llvm::createOpenMPOptCGSCCLegacyPass();
       (void) llvm::createArgumentPromotionPass();
       (void) llvm::createAlignmentFromAssumptionsPass();
       (void) llvm::createBasicAAWrapperPass();
       (void) llvm::createSCEVAAWrapperPass();
       (void) llvm::createTypeBasedAAWrapperPass();
       (void) llvm::createScopedNoAliasAAWrapperPass();
       (void) llvm::createBoundsCheckingLegacyPass();
       (void) llvm::createBreakCriticalEdgesPass();
ForcePassLinking() {
       (void) llvm::createSpeculativeExecutionIfHasBranchDivergencePass();
       (void) llvm::createRewriteSymbolsPass();
       (void) llvm::createStraightLineStrengthReducePass();
       (void) llvm::createMemDerefPrinter();
       (void) llvm::createMustExecutePrinter();
       (void) llvm::createMustBeExecutedContextPrinter();
       (void) llvm::createFloat2IntPass();
       (void) llvm::createEliminateAvailableExternallyPass();
-      (void)llvm::createScalarizeMaskedMemIntrinLegacyPass();
+      (void) llvm::createScalarizeMaskedMemIntrinLegacyPass();
       (void) llvm::createWarnMissedTransformationsPass();
       (void) llvm::createHardwareLoopsPass();
       (void) llvm::createInjectTLIMappingsLegacyPass();
       (void) llvm::createUnifyLoopExitsPass();
       (void) llvm::createFixIrreduciblePass();
 
       (void)new llvm::IntervalPartition();
       (void)new llvm::ScalarEvolutionWrapperPass();

llvm/include/llvm/Transforms/IPO.h

 /// createArgumentPromotionPass - This pass promotes "by reference" arguments to
 /// be passed by value if the number of elements passed is smaller or
 /// equal to maxElements (maxElements == 0 means always promote).
 ///
 Pass *createArgumentPromotionPass(unsigned maxElements = 3);
 
 //===----------------------------------------------------------------------===//
 /// createOpenMPOptLegacyPass - OpenMP specific optimizations.
-Pass *createOpenMPOptLegacyPass();
+Pass *createOpenMPOptCGSCCLegacyPass();
 
 //===----------------------------------------------------------------------===//
 /// createIPSCCPPass - This pass propagates constants from call sites into the
 /// bodies of functions, and keeps track of whether basic blocks are executable
 /// in the process.
 ///
 ModulePass *createIPSCCPPass();
 

llvm/include/llvm/Transforms/IPO/OpenMPOpt.h

 } // namespace omp
 
 /// OpenMP optimizations pass.
 class OpenMPOptPass : public PassInfoMixin<OpenMPOptPass> {
   /// Helper to remember if the module contains OpenMP (runtime calls).
   omp::OpenMPInModule OMPInModule;
 
 public:
+  PreservedAnalyses run(Module &M, ModuleAnalysisManager &AM);
+};
+
+class OpenMPOptCGSCCPass : public PassInfoMixin<OpenMPOptCGSCCPass> {
+  /// Helper to remember if the module contains OpenMP (runtime calls).
+  omp::OpenMPInModule OMPInModule;
+
+public:
   PreservedAnalyses run(LazyCallGraph::SCC &C, CGSCCAnalysisManager &AM,
                         LazyCallGraph &CG, CGSCCUpdateResult &UR);
 };
 
 } // end namespace llvm
 
 #endif // LLVM_TRANSFORMS_IPO_OPENMPOPT_H

llvm/lib/Passes/PassBuilder.cpp

PassBuilder::buildInlinerPipeline(OptimizationLevel Level,
   // When at O3 add argument promotion to the pass pipeline.
   // FIXME: It isn't at all clear why this should be limited to O3.
   if (Level == OptimizationLevel::O3)
     MainCGPipeline.addPass(ArgumentPromotionPass());
 
   // Try to perform OpenMP specific optimizations. This is a (quick!) no-op if
   // there are no OpenMP runtime calls present in the module.
   if (Level == OptimizationLevel::O2 || Level == OptimizationLevel::O3)
-    MainCGPipeline.addPass(OpenMPOptPass());
+    MainCGPipeline.addPass(OpenMPOptCGSCCPass());
 
   for (auto &C : CGSCCOptimizerLateEPCallbacks)
     C(MainCGPipeline, Level);
 
   // Lastly, add the core function simplification pipeline nested inside the
   // CGSCC walk.
   MainCGPipeline.addPass(createCGSCCToFunctionPassAdaptor(
       buildFunctionSimplificationPipeline(Level, Phase)));
if (Phase != ThinOrFullLTOPhase::ThinLTOPreLink &&
       // We perform early indirect call promotion here, before globalopt.
       // This is important for the ThinLTO backend phase because otherwise
       // imported available_externally functions look unreferenced and are
       // removed.
       MPM.addPass(
           PGOIndirectCallPromotion(true /* IsInLTO */, true /* SamplePGO */));
   }
 
+  // Try to perform OpenMP specific optimizations on the module. This is a
+  // (quick!) no-op if there are no OpenMP runtime calls present in the module.
+  if (Level == OptimizationLevel::O2 || Level == OptimizationLevel::O3)
+    MPM.addPass(OpenMPOptPass());
+
   if (AttributorRun & AttributorRunOption::MODULE)
     MPM.addPass(AttributorPass());
 
   // Lower type metadata and the type.test intrinsic in the ThinLTO
   // post link pipeline after ICP. This is to enable usage of the type
   // tests in ICP sequences.
   if (Phase == ThinOrFullLTOPhase::ThinLTOPostLink)
     MPM.addPass(LowerTypeTestsPass(nullptr, nullptr, true));

llvm/lib/Passes/PassRegistry.def

 #undef MODULE_ANALYSIS
 
 #ifndef MODULE_PASS
 #define MODULE_PASS(NAME, CREATE_PASS)
 #endif
 MODULE_PASS("always-inline", AlwaysInlinerPass())
 MODULE_PASS("attributor", AttributorPass())
 MODULE_PASS("annotation2metadata", Annotation2MetadataPass())
+MODULE_PASS("openmp-opt", OpenMPOptPass())
 MODULE_PASS("called-value-propagation", CalledValuePropagationPass())
 MODULE_PASS("canonicalize-aliases", CanonicalizeAliasesPass())
 MODULE_PASS("cg-profile", CGProfilePass())
 MODULE_PASS("constmerge", ConstantMergePass())
 MODULE_PASS("cross-dso-cfi", CrossDSOCFIPass())
 MODULE_PASS("deadargelim", DeadArgumentEliminationPass())
 MODULE_PASS("elim-avail-extern", EliminateAvailableExternallyPass())
 MODULE_PASS("extract-blocks", BlockExtractorPass())

 #ifndef CGSCC_PASS
 #define CGSCC_PASS(NAME, CREATE_PASS)
 #endif
 CGSCC_PASS("argpromotion", ArgumentPromotionPass())
 CGSCC_PASS("invalidate<all>", InvalidateAllAnalysesPass())
 CGSCC_PASS("function-attrs", PostOrderFunctionAttrsPass())
 CGSCC_PASS("attributor-cgscc", AttributorCGSCCPass())
 CGSCC_PASS("inline", InlinerPass())
-CGSCC_PASS("openmpopt", OpenMPOptPass())
+CGSCC_PASS("openmp-opt-cgscc", OpenMPOptCGSCCPass())
 CGSCC_PASS("coro-split", CoroSplitPass())
 CGSCC_PASS("no-op-cgscc", NoOpCGSCCPass())
 #undef CGSCC_PASS
 
 #ifndef FUNCTION_ANALYSIS
 #define FUNCTION_ANALYSIS(NAME, CREATE_PASS)
 #endif
 FUNCTION_ANALYSIS("aa", AAManager())

llvm/lib/Transforms/IPO/IPO.cpp

 #include "llvm/InitializePasses.h"
 #include "llvm/Transforms/IPO.h"
 #include "llvm/Transforms/IPO/AlwaysInliner.h"
 #include "llvm/Transforms/IPO/FunctionAttrs.h"
 
 using namespace llvm;
 
 void llvm::initializeIPO(PassRegistry &Registry) {
-  initializeOpenMPOptLegacyPassPass(Registry);
+  initializeOpenMPOptCGSCCLegacyPassPass(Registry);
   initializeArgPromotionPass(Registry);
   initializeAnnotation2MetadataLegacyPass(Registry);
   initializeCalledValuePropagationLegacyPassPass(Registry);
   initializeConstantMergeLegacyPassPass(Registry);
   initializeCrossDSOCFIPass(Registry);
   initializeDAEPass(Registry);
   initializeDAHPass(Registry);
   initializeForceFunctionAttrsLegacyPassPass(Registry);

llvm/lib/Transforms/IPO/OpenMPOpt.cpp

struct OpenMPOpt {
 
   /// Check if any remarks are enabled for openmp-opt
   bool remarksEnabled() {
     auto &Ctx = M.getContext();
     return Ctx.getDiagHandlerPtr()->isAnyRemarkEnabled(DEBUG_TYPE);
   }
 
   /// Run all OpenMP optimizations on the underlying SCC/ModuleSlice.
-  bool run() {
+  bool run(bool IsModulePass) {
     if (SCC.empty())
       return false;
 
     bool Changed = false;
 
     LLVM_DEBUG(dbgs() << TAG << "Run on SCC with " << SCC.size()
                       << " functions in a slice with "
                       << OMPInfoCache.ModuleSlice.size() << " functions\n");
 
     if (PrintICVValues)
       printICVs();
     if (PrintOpenMPKernels)
       printKernels();
 
     Changed |= rewriteDeviceCodeStateMachine();
 
     Changed |= runAttributor();
 
     // Recollect uses, in case Attributor deleted any.
     OMPInfoCache.recollectUses();
 
     Changed |= deleteParallelRegions();
     if (HideMemoryTransferLatency)
       Changed |= hideMemTransfersLatency();
-    if (remarksEnabled())
+    if (remarksEnabled() && IsModulePass)
       analysisGlobalization();
     Changed |= deduplicateRuntimeCalls();
     if (EnableParallelRegionMerging) {
       if (mergeParallelRegions()) {
         deduplicateRuntimeCalls();
         Changed = true;
       }
     }
for (auto &It : BB2PRMap) {
         MergableCIs.clear();
       }
 
       if (!MergableCIsVector.empty()) {
         Changed = true;
 
         for (auto &MergableCIs : MergableCIsVector)
           Merge(MergableCIs, BB);
+        MergableCIsVector.clear();
       }
     }
 
     if (Changed) {
       /// Re-collect use for fork calls, emitted barrier calls, and
       /// any emitted master/end_master calls.
       OMPInfoCache.recollectUsesForFunction(OMPRTL___kmpc_fork_call);
       OMPInfoCache.recollectUsesForFunction(OMPRTL___kmpc_barrier);
AAICVTracker &AAICVTracker::createForPosition(const IRPosition &IRP,
   case IRPosition::IRP_FUNCTION:
     AA = new (A.Allocator) AAICVTrackerFunction(IRP, A);
     break;
   }
 
   return *AA;
 }
 
-PreservedAnalyses OpenMPOptPass::run(LazyCallGraph::SCC &C,
+PreservedAnalyses OpenMPOptPass::run(Module &M, ModuleAnalysisManager &AM) {
+  if (!containsOpenMP(M, OMPInModule))
+    return PreservedAnalyses::all();
+
+  if (DisableOpenMPOptimizations)
+    return PreservedAnalyses::all();
+
+  // Look at every function definition in the Module.
+  SmallVector<Function *, 16> SCC;
+  for (Function &Fn : M)
+    if (!Fn.isDeclaration())
+      SCC.push_back(&Fn);
+
+  if (SCC.empty())
+    return PreservedAnalyses::all();
+
+  FunctionAnalysisManager &FAM =
+      AM.getResult<FunctionAnalysisManagerModuleProxy>(M).getManager();
+
+  AnalysisGetter AG(FAM);
+
+  auto OREGetter = [&FAM](Function *F) -> OptimizationRemarkEmitter & {
+    return FAM.getResult<OptimizationRemarkEmitterAnalysis>(*F);
+  };
+
+  BumpPtrAllocator Allocator;
+  CallGraphUpdater CGUpdater;
+
+  SetVector<Function *> Functions(SCC.begin(), SCC.end());
+  OMPInformationCache InfoCache(M, AG, Allocator, /*CGSCC*/ Functions,
+                                OMPInModule.getKernels());
+
+  Attributor A(Functions, InfoCache, CGUpdater);
+
+  OpenMPOpt OMPOpt(SCC, CGUpdater, OREGetter, InfoCache, A);
+  bool Changed = OMPOpt.run(true);
+  if (Changed)
+    return PreservedAnalyses::none();
+
+  return PreservedAnalyses::all();
+}
+
+PreservedAnalyses OpenMPOptCGSCCPass::run(LazyCallGraph::SCC &C,
                                           CGSCCAnalysisManager &AM,
-                                     LazyCallGraph &CG, CGSCCUpdateResult &UR) {
+                                          LazyCallGraph &CG,
+                                          CGSCCUpdateResult &UR) {
   if (!containsOpenMP(*C.begin()->getFunction().getParent(), OMPInModule))
     return PreservedAnalyses::all();
 
   if (DisableOpenMPOptimizations)
     return PreservedAnalyses::all();
 
   SmallVector<Function *, 16> SCC;
   // If there are kernels in the module, we have to run on all SCC's.
FunctionAnalysisManager &FAM =
       AM.getResult<FunctionAnalysisManagerCGSCCProxy>(C, CG).getManager();
 
   AnalysisGetter AG(FAM);
 
   auto OREGetter = [&FAM](Function *F) -> OptimizationRemarkEmitter & {
     return FAM.getResult<OptimizationRemarkEmitterAnalysis>(*F);
   };
 
+  BumpPtrAllocator Allocator;
   CallGraphUpdater CGUpdater;
   CGUpdater.initialize(CG, C, AM, UR);
 
   SetVector<Function *> Functions(SCC.begin(), SCC.end());
-  BumpPtrAllocator Allocator;
   OMPInformationCache InfoCache(*(Functions.back()->getParent()), AG, Allocator,
                                 /*CGSCC*/ Functions, OMPInModule.getKernels());
 
   Attributor A(Functions, InfoCache, CGUpdater);
 
   OpenMPOpt OMPOpt(SCC, CGUpdater, OREGetter, InfoCache, A);
-  bool Changed = OMPOpt.run();
+  bool Changed = OMPOpt.run(false);
   if (Changed)
     return PreservedAnalyses::none();
 
   return PreservedAnalyses::all();
 }
-
 namespace {
 
-struct OpenMPOptLegacyPass : public CallGraphSCCPass {
+struct OpenMPOptCGSCCLegacyPass : public CallGraphSCCPass {
   CallGraphUpdater CGUpdater;
   OpenMPInModule OMPInModule;
   static char ID;
 
-  OpenMPOptLegacyPass() : CallGraphSCCPass(ID) {
-    initializeOpenMPOptLegacyPassPass(*PassRegistry::getPassRegistry());
+  OpenMPOptCGSCCLegacyPass() : CallGraphSCCPass(ID) {
+    initializeOpenMPOptCGSCCLegacyPassPass(*PassRegistry::getPassRegistry());
   }
 
   void getAnalysisUsage(AnalysisUsage &AU) const override {
     CallGraphSCCPass::getAnalysisUsage(AU);
   }
 
   bool doInitialization(CallGraph &CG) override {
     // Disable the pass if there is no OpenMP (runtime call) in the module.
bool runOnSCC(CallGraphSCC &CGSCC) override {
     BumpPtrAllocator Allocator;
     OMPInformationCache InfoCache(
         *(Functions.back()->getParent()), AG, Allocator,
         /*CGSCC*/ Functions, OMPInModule.getKernels());
 
     Attributor A(Functions, InfoCache, CGUpdater);
 
     OpenMPOpt OMPOpt(SCC, CGUpdater, OREGetter, InfoCache, A);
-    return OMPOpt.run();
+    return OMPOpt.run(false);
   }
 
   bool doFinalization(CallGraph &CG) override { return CGUpdater.finalize(); }
 };
 
 } // end anonymous namespace
 
 void OpenMPInModule::identifyKernels(Module &M) {
#include "llvm/Frontend/OpenMP/OMPKinds.def"
   if (OMPInModule.isKnown() && OMPInModule) {
     OMPInModule.identifyKernels(M);
     return true;
   }
 
   return OMPInModule = false;
 }
 
-char OpenMPOptLegacyPass::ID = 0;
+char OpenMPOptCGSCCLegacyPass::ID = 0;
 
-INITIALIZE_PASS_BEGIN(OpenMPOptLegacyPass, "openmpopt",
+INITIALIZE_PASS_BEGIN(OpenMPOptCGSCCLegacyPass, "openmp-opt-cgscc",
                       "OpenMP specific optimizations", false, false)
 INITIALIZE_PASS_DEPENDENCY(CallGraphWrapperPass)
-INITIALIZE_PASS_END(OpenMPOptLegacyPass, "openmpopt",
+INITIALIZE_PASS_END(OpenMPOptCGSCCLegacyPass, "openmp-opt-cgscc",
                     "OpenMP specific optimizations", false, false)
 
-Pass *llvm::createOpenMPOptLegacyPass() { return new OpenMPOptLegacyPass(); }
+Pass *llvm::createOpenMPOptCGSCCLegacyPass() {
+  return new OpenMPOptCGSCCLegacyPass();
+}

llvm/lib/Transforms/IPO/PassManagerBuilder.cpp

void PassManagerBuilder::populateModulePassManager(
 
   // Infer attributes on declarations, call sites, arguments, etc. for an SCC.
   if (AttributorRun & AttributorRunOption::CGSCC)
     MPM.add(createAttributorCGSCCLegacyPass());
 
   // Try to perform OpenMP specific optimizations. This is a (quick!) no-op if
   // there are no OpenMP runtime calls present in the module.
   if (OptLevel > 1)
-    MPM.add(createOpenMPOptLegacyPass());
+    MPM.add(createOpenMPOptCGSCCLegacyPass());
 
   MPM.add(createPostOrderFunctionAttrsLegacyPass());
   if (OptLevel > 2)
     MPM.add(createArgumentPromotionPass()); // Scalarize uninlined fn args
 
   addExtensionsToPM(EP_CGSCCOptimizerLate, MPM);
   addFunctionSimplificationPasses(MPM);
 
void PassManagerBuilder::addLTOOptimizationPasses(legacy::PassManagerBase &PM) {
 
   // Infer attributes on declarations, call sites, arguments, etc. for an SCC.
   if (AttributorRun & AttributorRunOption::CGSCC)
     PM.add(createAttributorCGSCCLegacyPass());
 
   // Try to perform OpenMP specific optimizations. This is a (quick!) no-op if
   // there are no OpenMP runtime calls present in the module.
   if (OptLevel > 1)
-    PM.add(createOpenMPOptLegacyPass());
+    PM.add(createOpenMPOptCGSCCLegacyPass());
 
   // Optimize globals again if we ran the inliner.
   if (RunInliner)
     PM.add(createGlobalOptimizerPass());
   PM.add(createGlobalDCEPass()); // Remove dead functions.
 
   // If we didn't decide to inline a function, check to see if we can
   // transform it to pass arguments by value instead of by reference.

llvm/test/Transforms/OpenMP/add_attributes.ll

-; RUN: opt < %s -S -openmpopt        | FileCheck %s
-; RUN: opt < %s -S -passes=openmpopt | FileCheck %s
-; RUN: opt < %s -S -openmpopt        -openmp-ir-builder-optimistic-attributes | FileCheck %s --check-prefix=OPTIMISTIC
-; RUN: opt < %s -S -passes=openmpopt -openmp-ir-builder-optimistic-attributes | FileCheck %s --check-prefix=OPTIMISTIC
+; RUN: opt < %s -S -openmp-opt-cgscc        | FileCheck %s
+; RUN: opt < %s -S -passes=openmp-opt-cgscc | FileCheck %s
+; RUN: opt < %s -S -openmp-opt-cgscc        -openmp-ir-builder-optimistic-attributes | FileCheck %s --check-prefix=OPTIMISTIC
+; RUN: opt < %s -S -passes=openmp-opt-cgscc -openmp-ir-builder-optimistic-attributes | FileCheck %s --check-prefix=OPTIMISTIC
 target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128"
 
 %struct.omp_lock_t = type { i8* }
 %struct.omp_nest_lock_t = type { i8* }
 %struct.ident_t = type { i32, i32, i32, i32, i8* }
 
 define void @call_all(i32 %schedule, %struct.omp_lock_t* %lock, i32 %lock_hint, %struct.omp_nest_lock_t* %nest_lock, i32 %i, i8* %s, i64 %st, i8* %vp, double %d, i32 %proc_bind, i64 %allocator_handle, i8* %cp, i64 %event_handle, i32 %pause_resource) {
 entry:

llvm/test/Transforms/OpenMP/add_attributes_amdgcn.ll

-; RUN: opt < %s -S -openmpopt        | FileCheck %s
-; RUN: opt < %s -S -passes=openmpopt | FileCheck %s
-; RUN: opt < %s -S -openmpopt        -openmp-ir-builder-optimistic-attributes | FileCheck %s --check-prefix=OPTIMISTIC
-; RUN: opt < %s -S -passes=openmpopt -openmp-ir-builder-optimistic-attributes | FileCheck %s --check-prefix=OPTIMISTIC
+; RUN: opt < %s -S -openmp-opt-cgscc        | FileCheck %s
+; RUN: opt < %s -S -passes=openmp-opt-cgscc | FileCheck %s
+; RUN: opt < %s -S -openmp-opt-cgscc        -openmp-ir-builder-optimistic-attributes | FileCheck %s --check-prefix=OPTIMISTIC
+; RUN: opt < %s -S -passes=openmp-opt-cgscc -openmp-ir-builder-optimistic-attributes | FileCheck %s --check-prefix=OPTIMISTIC
 
 target triple = "amdgcn-amd-amdhsa"
 
 define void @call_all(i64 %arg) {
   call void @__kmpc_syncwarp(i64 %arg)
   call i64 @__kmpc_warp_active_thread_mask()
   ret void
 }

llvm/test/Transforms/OpenMP/attributor_module_slice_reproducer.ll

 ; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --function-signature
-; RUN: opt -S -openmpopt < %s | FileCheck %s
-; RUN: opt -S -passes=openmpopt < %s | FileCheck %s
+; RUN: opt -S -openmp-opt-cgscc < %s | FileCheck %s
+; RUN: opt -S -passes=openmp-opt-cgscc < %s | FileCheck %s
 
 define internal fastcc void @"_omp$reduction$reduction_func14"() unnamed_addr {
   %call = call i8 @_ZStplIdESt7complexIT_ERKS2_S4_()
   ret void
 }
 
 define linkonce_odr hidden i8 @_ZStplIdESt7complexIT_ERKS2_S4_() local_unnamed_addr {
 ; CHECK-LABEL: define {{[^@]+}}@_ZStplIdESt7complexIT_ERKS2_S4_() local_unnamed_addr {

llvm/test/Transforms/OpenMP/dead_use.ll

 ; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --function-signature
-; RUN: opt -S -openmpopt < %s | FileCheck %s
-; RUN: opt -S -passes=openmpopt < %s | FileCheck %s
+; RUN: opt -S -openmp-opt-cgscc < %s | FileCheck %s
+; RUN: opt -S -passes=openmp-opt-cgscc < %s | FileCheck %s
 %struct.ident_t = type { i32, i32, i32, i32, i8* }
 
 @.str = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00", align 1
 @0 = private unnamed_addr global %struct.ident_t { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* @.str, i32 0, i32 0) }, align 8
 
 ; Function Attrs: nounwind uwtable
 define dso_local i32 @b() #0 {
 ; CHECK-LABEL: define {{[^@]+}}@b() #0

llvm/test/Transforms/OpenMP/deduplication.ll

 ; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --function-signature --scrub-attributes
-; RUN: opt -openmpopt -S < %s | FileCheck %s
-; RUN: opt -passes=openmpopt -S < %s | FileCheck %s
+; RUN: opt -openmp-opt-cgscc -S < %s | FileCheck %s
+; RUN: opt -passes=openmp-opt-cgscc -S < %s | FileCheck %s
 target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128"
 
 %struct.ident_t = type { i32, i32, i32, i32, i8* }
 
 @0 = private unnamed_addr constant %struct.ident_t { i32 0, i32 34, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* @.str0, i32 0, i32 0) }, align 8
 @1 = private unnamed_addr constant %struct.ident_t { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* @.str1, i32 0, i32 0) }, align 8
 @2 = private unnamed_addr constant %struct.ident_t { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* @.str2, i32 0, i32 0) }, align 8
 @.str0 = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00", align 1

llvm/test/Transforms/OpenMP/deduplication_remarks.ll

-; RUN: opt -openmpopt -pass-remarks=openmp-opt -disable-output < %s 2>&1 | FileCheck %s
-; RUN: opt -passes=openmpopt -pass-remarks=openmp-opt -disable-output < %s 2>&1 | FileCheck %s
+; RUN: opt -openmp-opt-cgscc -pass-remarks=openmp-opt -disable-output < %s 2>&1 | FileCheck %s
+; RUN: opt -passes=openmp-opt-cgscc -pass-remarks=openmp-opt -disable-output < %s 2>&1 | FileCheck %s
 ; ModuleID = 'deduplication_remarks.c'
 source_filename = "deduplication_remarks.c"
 target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128"
 target triple = "x86_64-pc-linux-gnu"
 
 %struct.ident_t = type { i32, i32, i32, i32, i8* }
 
 @0 = private unnamed_addr global %struct.ident_t { i32 0, i32 34, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* @.str0, i32 0, i32 0) }, align 8

llvm/test/Transforms/OpenMP/globalization_remarks.ll

-; RUN: opt -passes=openmpopt -pass-remarks-analysis=openmp-opt -disable-output < %s 2>&1 | FileCheck %s
+; RUN: opt -passes=openmp-opt -pass-remarks-analysis=openmp-opt -disable-output < %s 2>&1 | FileCheck %s
 ; ModuleID = 'declare_target_codegen_globalization.cpp'
 source_filename = "declare_target_codegen_globalization.cpp"
 target datalayout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64"
 target triple = "nvptx64-nvidia-cuda"
 
 %struct.ident_t = type { i32, i32, i32, i32, i8* }
 %struct._globalized_locals_ty = type { [32 x i32] }
 

llvm/test/Transforms/OpenMP/gpu_kernel_detection_remarks.ll

-; RUN: opt -passes=openmpopt -pass-remarks=openmp-opt -openmp-print-gpu-kernels -disable-output < %s 2>&1 | FileCheck %s --implicit-check-not=non_kernel
-; RUN: opt        -openmpopt -pass-remarks=openmp-opt -openmp-print-gpu-kernels -disable-output < %s 2>&1 | FileCheck %s --implicit-check-not=non_kernel
+; RUN: opt -passes=openmp-opt-cgscc -pass-remarks=openmp-opt -openmp-print-gpu-kernels -disable-output < %s 2>&1 | FileCheck %s --implicit-check-not=non_kernel
+; RUN: opt        -openmp-opt-cgscc -pass-remarks=openmp-opt -openmp-print-gpu-kernels -disable-output < %s 2>&1 | FileCheck %s --implicit-check-not=non_kernel
 
 ; CHECK-DAG: remark: <unknown>:0:0: OpenMP GPU kernel kernel1
 ; CHECK-DAG: remark: <unknown>:0:0: OpenMP GPU kernel kernel2
 
 define void @kernel1() {
   ret void
 }
 

llvm/test/Transforms/OpenMP/gpu_state_machine_function_ptr_replacement.ll

-; RUN: opt -S -passes=openmpopt -openmp-ir-builder-optimistic-attributes -pass-remarks=openmp-opt -openmp-print-gpu-kernels < %s | FileCheck %s
-; RUN: opt -S -passes=openmpopt -pass-remarks=openmp-opt -openmp-print-gpu-kernels < %s | FileCheck %s
-; RUN: opt -S        -openmpopt -pass-remarks=openmp-opt -openmp-print-gpu-kernels < %s | FileCheck %s
+; RUN: opt -S -passes=openmp-opt-cgscc -openmp-ir-builder-optimistic-attributes -pass-remarks=openmp-opt -openmp-print-gpu-kernels < %s | FileCheck %s
+; RUN: opt -S -passes=openmp-opt-cgscc -pass-remarks=openmp-opt -openmp-print-gpu-kernels < %s | FileCheck %s
+; RUN: opt -S        -openmp-opt-cgscc -pass-remarks=openmp-opt -openmp-print-gpu-kernels < %s | FileCheck %s
 
 ; C input used for this test:
 
 ; void bar(void) {
 ;     #pragma omp parallel
 ;     { }
 ; }
 ; void foo(void) {

llvm/test/Transforms/OpenMP/hide_mem_transfer_latency.ll

 ; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: -p --function-signature --scrub-attributes
-; RUN: opt -S -passes=openmpopt -aa-pipeline=basic-aa -openmp-hide-memory-transfer-latency < %s | FileCheck %s
+; RUN: opt -S -openmp-opt-cgscc -aa-pipeline=basic-aa -openmp-hide-memory-transfer-latency < %s | FileCheck %s
+; RUN: opt -S -passes=openmp-opt-cgscc -aa-pipeline=basic-aa -openmp-hide-memory-transfer-latency < %s | FileCheck %s
 target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128"
 
 ; CHECK: %struct.__tgt_async_info = type { i8* }
 
 %struct.ident_t = type { i32, i32, i32, i32, i8* }
 %struct.__tgt_offload_entry = type { i8*, i8*, i64, i32, i32 }
 
 @.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 35]

llvm/test/Transforms/OpenMP/icv_remarks.ll

-; RUN: opt -passes=openmpopt -pass-remarks=openmp-opt -openmp-print-icv-values -disable-output < %s 2>&1 | FileCheck %s
-; RUN: opt -openmpopt -pass-remarks=openmp-opt -openmp-print-icv-values -disable-output < %s 2>&1 | FileCheck %s
+; RUN: opt -passes=openmp-opt-cgscc -pass-remarks=openmp-opt -openmp-print-icv-values -disable-output < %s 2>&1 | FileCheck %s
+; RUN: opt -openmp-opt-cgscc -pass-remarks=openmp-opt -openmp-print-icv-values -disable-output < %s 2>&1 | FileCheck %s
 
 ; ModuleID = 'icv_remarks.c'
 source_filename = "icv_remarks.c"
 target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128"
 target triple = "x86_64-unknown-linux-gnu"
 
 %struct.ident_t = type { i32, i32, i32, i32, i8* }
 

llvm/test/Transforms/OpenMP/icv_tracking.ll

 ; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --function-signature
-; RUN: opt -S -openmpopt < %s | FileCheck %s
-; RUN: opt -S -passes=openmpopt < %s | FileCheck %s
+; RUN: opt -S -openmp-opt-cgscc < %s | FileCheck %s
+; RUN: opt -S -passes=openmp-opt-cgscc < %s | FileCheck %s
 
 %struct.ident_t = type { i32, i32, i32, i32, i8* }
 
 @.str = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00", align 1
 @0 = private unnamed_addr global %struct.ident_t { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* @.str, i32 0, i32 0) }, align 8
 
 ; doesn't modify any ICVs.
 define i32 @icv_free_use(i32 %0) {

llvm/test/Transforms/OpenMP/parallel_deletion.ll

 ; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --function-signature
-; RUN: opt -S -attributor -openmpopt  < %s | FileCheck %s
-; RUN: opt -S -passes='attributor,cgscc(openmpopt)'  < %s | FileCheck %s
+; RUN: opt -S -attributor -openmp-opt-cgscc  < %s | FileCheck %s
+; RUN: opt -S -passes='attributor,cgscc(openmp-opt-cgscc)'  < %s | FileCheck %s
 ;
 target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128"
 
 %struct.ident_t = type { i32, i32, i32, i32, i8* }
 
 @.str = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00", align 1
 @0 = private unnamed_addr global %struct.ident_t { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* @.str, i32 0, i32 0) }, align 8
 @1 = private unnamed_addr global %struct.ident_t { i32 0, i32 322, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* @.str, i32 0, i32 0) }, align 8

llvm/test/Transforms/OpenMP/parallel_deletion_remarks.ll

-; RUN: opt -S -pass-remarks=openmp-opt -attributor -openmpopt -disable-output < %s 2>&1 | FileCheck %s
-; RUN: opt -S -pass-remarks=openmp-opt -passes='attributor,cgscc(openmpopt)' -disable-output < %s 2>&1 | FileCheck %s
+; RUN: opt -S -pass-remarks=openmp-opt -attributor -openmp-opt-cgscc -disable-output < %s 2>&1 | FileCheck %s
+; RUN: opt -S -pass-remarks=openmp-opt -passes='attributor,cgscc(openmp-opt-cgscc)' -disable-output < %s 2>&1 | FileCheck %s
 ; ModuleID = 'parallel_deletion_remarks.ll'
 source_filename = "parallel_deletion_remarks.c"
 target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128"
 target triple = "x86_64-pc-linux-gnu"
 
 %struct.ident_t = type { i32, i32, i32, i32, i8* }
 
 @.str = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00", align 1

llvm/test/Transforms/OpenMP/parallel_region_merging.ll

 ; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs
-; RUN: opt -S -attributor -openmpopt -openmp-opt-enable-merging -enable-new-pm=0 < %s | FileCheck %s
-; RUN: opt -S -aa-pipeline= -passes='attributor,cgscc(openmpopt)' -openmp-opt-enable-merging  < %s | FileCheck %s
+; RUN: opt -S -attributor -openmp-opt-cgscc -openmp-opt-enable-merging -enable-new-pm=0 < %s | FileCheck %s
+; RUN: opt -S -aa-pipeline= -passes='attributor,cgscc(openmp-opt-cgscc)' -openmp-opt-enable-merging  < %s | FileCheck %s
 ; #include <omp.h>
 ; void foo();
 ; void use(int);
 ; void usef(float);
 ; void merge(int a) {
 ; #pragma omp parallel
 ;     {
 ;         use(a);

llvm/test/Transforms/OpenMP/rtf_type_checking.ll

-; RUN: opt -S -openmpopt -stats < %s 2>&1 -enable-new-pm=0 | FileCheck %s --check-prefixes=CHECK,LPM
-; RUN: opt -S -passes='devirt<2>(cgscc(openmpopt))' -stats -debug-pass-manager < %s 2>&1 | FileCheck %s --check-prefixes=CHECK,NPM
-; RUN: opt -S -attributor -openmpopt -stats < %s 2>&1 -enable-new-pm=0 | FileCheck %s --check-prefixes=CHECK,LPM
-; RUN: opt -S -passes='attributor,cgscc(devirt<2>(openmpopt))' -stats -debug-pass-manager < %s 2>&1 | FileCheck %s --check-prefixes=CHECK,NPM
+; RUN: opt -S -openmp-opt-cgscc -stats < %s 2>&1 -enable-new-pm=0 | FileCheck %s --check-prefixes=CHECK,LPM
+; RUN: opt -S -passes='devirt<2>(cgscc(openmp-opt-cgscc))' -stats -debug-pass-manager < %s 2>&1 | FileCheck %s --check-prefixes=CHECK,NPM
+; RUN: opt -S -attributor -openmp-opt-cgscc -stats < %s 2>&1 -enable-new-pm=0 | FileCheck %s --check-prefixes=CHECK,LPM
+; RUN: opt -S -passes='attributor,cgscc(devirt<2>(openmp-opt-cgscc))' -stats -debug-pass-manager < %s 2>&1 | FileCheck %s --check-prefixes=CHECK,NPM
 ; REQUIRES: asserts
 
 target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128"
 
 %struct.ident_t = type { i32, i32, i32, i32, i8* }
 
 @.str = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00", align 1
 @0 = private unnamed_addr global %struct.ident_t { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* @.str, i32 0, i32 0) }, align 8

 !llvm.module.flags = !{!0}
 !llvm.ident = !{!1}
 
 !0 = !{i32 1, !"wchar_size", i32 4}
 !1 = !{!"clang"}
 !2 = !{!3}
 !3 = !{i64 2, i64 -1, i64 -1, i1 true}
 
-; NPM: Running pass: OpenMPOptPass on (.omp_outlined.)
-; NPM-NOT: Running pass: OpenMPOptPass on (.omp_outlined.)
-; NPM: Running pass: OpenMPOptPass on (main)
-; NPM-NOT: Running pass: OpenMPOptPass on (main)
+; NPM: Running pass: OpenMPOptCGSCCPass on (.omp_outlined.)
+; NPM-NOT: Running pass: OpenMPOptCGSCCPass on (.omp_outlined.)
+; NPM: Running pass: OpenMPOptCGSCCPass on (main)
+; NPM-NOT: Running pass: OpenMPOptCGSCCPass on (main)
 ; ===-------------------------------------------------------------------------===
 ;                         ... Statistics Collected ...
 ; ===-------------------------------------------------------------------------===
 ;
 ; LPM: 1 cgscc-passmgr - Maximum CGSCCPassMgr iterations on one SCC
 ; CHECK: 2 openmp-opt{{.*}}Number of OpenMP runtime functions identified
 ;
 ; There are two matches since the pass is run once per function.

llvm/test/Transforms/OpenMP/values_in_offload_arrays.ll

-; RUN: opt -S -passes=openmpopt -aa-pipeline=basic-aa -openmp-hide-memory-transfer-latency -debug-only=openmp-opt < %s 2>&1 | FileCheck %s
+; RUN: opt -S -openmp-opt-cgscc -aa-pipeline=basic-aa -openmp-hide-memory-transfer-latency -debug-only=openmp-opt < %s 2>&1 | FileCheck %s
+; RUN: opt -S -passes=openmp-opt-cgscc -aa-pipeline=basic-aa -openmp-hide-memory-transfer-latency -debug-only=openmp-opt < %s 2>&1 | FileCheck %s
 ; REQUIRES: asserts
 
 target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128"
 
 @.__omp_offloading_heavyComputation.region_id = weak constant i8 0
 @.offload_maptypes. = private unnamed_addr constant [2 x i64] [i64 35, i64 35]
 
 %struct.ident_t = type { i32, i32, i32, i32, i8* }

llvm/test/Transforms/PhaseOrdering/openmp-opt-module.ll

+; RUN: opt -O2 -pass-remarks-analysis=openmp-opt -enable-new-pm < %s 2>&1 | FileCheck %s --check-prefix=MODULE
+target datalayout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64"
+
+@.str = private unnamed_addr constant [13 x i8] c"Alloc Shared\00", align 1
+
+; MODULE: remark: openmp_opt_module.c:5:7: Found thread data sharing on the GPU. Expect degraded performance due to data globalization.
+
+define void @foo() {
+entry:
+  %x = call i8* @__kmpc_data_sharing_push_stack(i64 4, i16 0), !dbg !7
+  %x_on_stack = bitcast i8* %x to i32*
+  %0 = bitcast i32* %x_on_stack to i8*
+  call void @use(i8* %0)
+  call void @__kmpc_data_sharing_pop_stack(i8* %x)
+  ret void
+}
+
+define void @use(i8* %0) {
+entry:
+  %.addr = alloca i8*, align 8
+  store i8* %0, i8** %.addr, align 8
+  ret void
+}
+
+define internal i8* @__kmpc_data_sharing_push_stack(i64 %DataSize, i16 %shared) {
+entry:
+  %call = call i8* @_Z10SafeMallocmPKc(i64 %DataSize, i8* getelementptr inbounds ([13 x i8], [13 x i8]* @.str, i64 0, i64 0)) #11
+  ret i8* %call
+}
+
+; Function Attrs: convergent nounwind mustprogress
+declare i8* @_Z10SafeMallocmPKc(i64 %size, i8* nocapture readnone %msg)
+
+declare void @__kmpc_data_sharing_pop_stack(i8*)
+
+!llvm.dbg.cu = !{!0}
+!llvm.module.flags = !{!3, !4}
+
+!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang version 12.0.0", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, enums: !2, splitDebugInlining: false, nameTableKind: None)
+!1 = !DIFile(filename: "openmp_opt_module.c", directory: "/tmp/openmp_opt_module.c")
+!2 = !{}
+!3 = !{i32 2, !"Debug Info Version", i32 3}
+!4 = !{i32 1, !"wchar_size", i32 4}
+!5 = distinct !DISubprogram(name: "foo", scope: !1, file: !1, line: 1, type: !6, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !0, retainedNodes: !2)
+!6 = !DISubroutineType(types: !2)
+!7 = !DILocation(line: 5, column: 7, scope: !5)