[Diagnostics] Allow emitting analysis and missed remarks on functions

[jhuber6]

Summary

Currently, only OptimizationRemarks can be emitted using a Function.
Add constructors to allow this for OptimizationRemarksAnalysis and
OptimizationRemarkMissed as well.

[jhuber6]

Formatting.

[jdoerfert]

Diagnostic Info parts seems very mechanical and ok to me.
OpenMPOpt parts make sense.
Test seem to be always tied to an analysis (e.g., D29023) so LGTM

[thegameg]

LGTM on the new constructors.

[jhuber6]

Fixing Clang tests.

clang/test/OpenMP/remarks_parallel_in_multiple_target_state_machines.c

-// RUN: %clang_cc1                                 -verify=host                                                              -Rpass=openmp -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
-// RUN: %clang_cc1                                 -verify=all,safe                                                          -Rpass=openmp -fopenmp -O2 -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t.out
-// RUN: %clang_cc1 -fexperimental-new-pass-manager -verify=all,safe                                                          -Rpass=openmp -fopenmp -O2 -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t.out
+// RUN: %clang_cc1                                 -verify=host      -Rpass=openmp-opt -Rpass-analysis=openmp -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1                                 -verify=all,safe  -Rpass=openmp-opt -Rpass-analysis=openmp -fopenmp -O2 -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t.out
+// RUN: %clang_cc1 -fexperimental-new-pass-manager -verify=all,safe  -Rpass=openmp-opt -Rpass-analysis=openmp -fopenmp -O2 -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t.out
 
 // host-no-diagnostics
 
 void bar1(void) {
 #pragma omp parallel // #0
                      // all-remark@#0 {{Found a parallel region that is called in a target region but not part of a combined target construct nor nested inside a target construct without intermediate code. This can lead to excessive register usage for unrelated target regions in the same translation unit due to spurious call edges assumed by ptxas.}}
                      // safe-remark@#0 {{Parallel region is used in unexpected ways; will not attempt to rewrite the state machine.}}
                      // force-remark@#0 {{Specialize parallel region that is only reached from a single target region to avoid spurious call edges and excessive register usage in other target regions. (parallel region ID: __omp_outlined__2_wrapper, kernel ID: <NONE>}}
#pragma omp parallel
   {
   }
 
 #pragma omp target teams distribute parallel for
   for (int i = 0; i < 100; ++i) {
   }
 }
 
-// all-remark@* 5 {{OpenMP runtime call __kmpc_global_thread_num moved to}}
+// all-remark@* 5 {{OpenMP runtime call __kmpc_global_thread_num moved to beginning of OpenMP region}}
 // all-remark@* 12 {{OpenMP runtime call __kmpc_global_thread_num deduplicated}}

clang/test/OpenMP/remarks_parallel_in_target_state_machine.c

-// RUN: %clang_cc1                                 -verify=host -Rpass=openmp -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
-// RUN: %clang_cc1                                 -verify      -Rpass=openmp -fopenmp -O2 -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t.out
-// RUN: %clang_cc1 -fexperimental-new-pass-manager -verify      -Rpass=openmp -fopenmp -O2 -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t.out
+// RUN: %clang_cc1                                 -verify=host -Rpass=openmp -Rpass-analysis=openmp-opt -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1                                 -verify      -Rpass=openmp -Rpass-analysis=openmp-opt -fopenmp -O2 -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t.out
+// RUN: %clang_cc1 -fexperimental-new-pass-manager -verify      -Rpass=openmp -Rpass-analysis=openmp-opt -fopenmp -O2 -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t.out
 
 // host-no-diagnostics
 
 void bar(void) {
 #pragma omp parallel // #1                                                                                                                                                                                                                                                                                                                                           \
                      // expected-remark@#1 {{Found a parallel region that is called in a target region but not part of a combined target construct nor nested inside a target construct without intermediate code. This can lead to excessive register usage for unrelated target regions in the same translation unit due to spurious call edges assumed by ptxas.}} \
                      // expected-remark@#1 {{Parallel region is used in unexpected ways; will not attempt to rewrite the state machine.}}
   {
#pragma omp parallel
   {
   }
 
 #pragma omp target teams distribute parallel for
   for (int i = 0; i < 100; ++i) {
   }
 }
 
-// expected-remark@* {{OpenMP runtime call __kmpc_global_thread_num moved to}}
+// expected-remark@* {{OpenMP runtime call __kmpc_global_thread_num moved to beginning of OpenMP region}}
 // expected-remark@* 2 {{OpenMP runtime call __kmpc_global_thread_num deduplicated}}

llvm/include/llvm/IR/DiagnosticInfo.h

OptimizationRemarkMissed(const char *PassName, StringRef RemarkName,
                            const DiagnosticLocation &Loc,
                            const Value *CodeRegion);
 
   /// Same as above but \p Inst is used to derive code region and debug
   /// location.
   OptimizationRemarkMissed(const char *PassName, StringRef RemarkName,
                            const Instruction *Inst);
 
+  /// Same as above but \p F is used to derive code region and debug
+  /// location.
+  OptimizationRemarkMissed(const char *PassName, StringRef RemarkName,
+                           const Function *F);
+
   static bool classof(const DiagnosticInfo *DI) {
     return DI->getKind() == DK_OptimizationRemarkMissed;
   }
 
   /// \see DiagnosticInfoOptimizationBase::isEnabled.
   bool isEnabled() const override;
 
 private:
OptimizationRemarkAnalysis(const char *PassName, StringRef Prepend,
                              const OptimizationRemarkAnalysis &Orig)
       : DiagnosticInfoIROptimization(PassName, Prepend, Orig) {}
 
   /// Same as above but \p Inst is used to derive code region and debug
   /// location.
   OptimizationRemarkAnalysis(const char *PassName, StringRef RemarkName,
                              const Instruction *Inst);
 
+  /// Same as above but \p F is used to derive code region and debug
+  /// location.
+  OptimizationRemarkAnalysis(const char *PassName, StringRef RemarkName,
+                             const Function *F);
+
   static bool classof(const DiagnosticInfo *DI) {
     return DI->getKind() == DK_OptimizationRemarkAnalysis;
   }
 
   /// \see DiagnosticInfoOptimizationBase::isEnabled.
   bool isEnabled() const override;
 
   static const char *AlwaysPrint;

llvm/lib/IR/DiagnosticInfo.cpp

 OptimizationRemarkMissed::OptimizationRemarkMissed(const char *PassName,
                                                    StringRef RemarkName,
                                                    const Instruction *Inst)
     : DiagnosticInfoIROptimization(DK_OptimizationRemarkMissed, DS_Remark,
                                    PassName, RemarkName,
                                    *Inst->getParent()->getParent(),
                                    Inst->getDebugLoc(), Inst->getParent()) {}
 
+OptimizationRemarkMissed::OptimizationRemarkMissed(const char *PassName,
+                                                   StringRef RemarkName,
+                                                   const Function *Func)
+    : DiagnosticInfoIROptimization(
+          DK_OptimizationRemarkMissed, DS_Remark, PassName, RemarkName, *Func,
+          Func->getSubprogram(), getFirstFunctionBlock(Func)) {}
+
 bool OptimizationRemarkMissed::isEnabled() const {
   const Function &Fn = getFunction();
   LLVMContext &Ctx = Fn.getContext();
   return Ctx.getDiagHandlerPtr()->isMissedOptRemarkEnabled(getPassName());
 }
 
 OptimizationRemarkAnalysis::OptimizationRemarkAnalysis(
     const char *PassName, StringRef RemarkName, const DiagnosticLocation &Loc,

 
 OptimizationRemarkAnalysis::OptimizationRemarkAnalysis(
     enum DiagnosticKind Kind, const char *PassName, StringRef RemarkName,
     const DiagnosticLocation &Loc, const Value *CodeRegion)
     : DiagnosticInfoIROptimization(Kind, DS_Remark, PassName, RemarkName,
                                    *cast<BasicBlock>(CodeRegion)->getParent(),
                                    Loc, CodeRegion) {}
 
+OptimizationRemarkAnalysis::OptimizationRemarkAnalysis(const char *PassName,
+                                                       StringRef RemarkName,
+                                                       const Function *Func)
+    : DiagnosticInfoIROptimization(
+          DK_OptimizationRemarkAnalysis, DS_Remark, PassName, RemarkName, *Func,
+          Func->getSubprogram(), getFirstFunctionBlock(Func)) {}
+
 bool OptimizationRemarkAnalysis::isEnabled() const {
   const Function &Fn = getFunction();
   LLVMContext &Ctx = Fn.getContext();
   return Ctx.getDiagHandlerPtr()->isAnalysisRemarkEnabled(getPassName()) ||
          shouldAlwaysPrint();
 }
 
 void DiagnosticInfoMIRParser::print(DiagnosticPrinter &DP) const {

llvm/lib/Transforms/IPO/OpenMPOpt.cpp

struct OpenMPOpt {
   /// FIXME: This should be done from the Attributor once it is added.
   void printICVs() const {
     InternalControlVar ICVs[] = {ICV_nthreads, ICV_active_levels, ICV_cancel,
                                  ICV_proc_bind};
 
     for (Function *F : OMPInfoCache.ModuleSlice) {
       for (auto ICV : ICVs) {
         auto ICVInfo = OMPInfoCache.ICVs[ICV];
-        auto Remark = [&](OptimizationRemark OR) {
-          return OR << "OpenMP ICV " << ore::NV("OpenMPICV", ICVInfo.Name)
+        auto Remark = [&](OptimizationRemarkAnalysis ORA) {
+          return ORA << "OpenMP ICV " << ore::NV("OpenMPICV", ICVInfo.Name)
                      << " Value: "
                      << (ICVInfo.InitValue
                              ? ICVInfo.InitValue->getValue().toString(10, true)
                              : "IMPLEMENTATION_DEFINED");
         };
 
-        emitRemarkOnFunction(F, "OpenMPICVTracker", Remark);
+        emitRemark<OptimizationRemarkAnalysis>(F, "OpenMPICVTracker", Remark);
       }
     }
   }
 
   /// Print OpenMP GPU kernels for testing.
   void printKernels() const {
     for (Function *F : SCC) {
       if (!OMPInfoCache.Kernels.count(F))
         continue;
 
-      auto Remark = [&](OptimizationRemark OR) {
-        return OR << "OpenMP GPU kernel "
+      auto Remark = [&](OptimizationRemarkAnalysis ORA) {
+        return ORA << "OpenMP GPU kernel "
                    << ore::NV("OpenMPGPUKernel", F->getName()) << "\n";
       };
 
-      emitRemarkOnFunction(F, "OpenMPGPU", Remark);
+      emitRemark<OptimizationRemarkAnalysis>(F, "OpenMPGPU", Remark);
     }
   }
 
   /// Return the call if \p U is a callee use in a regular call. If \p RFI is
   /// given it has to be the callee or a nullptr is returned.
   static CallInst *getCallIfRegularCall(
       Use &U, OMPInformationCache::RuntimeFunctionInfo *RFI = nullptr) {
     CallInst *CI = dyn_cast<CallInst>(U.getUser());
bool deduplicateRuntimeCalls(Function &F,
 
     if (!ReplVal) {
       for (Use *U : *UV)
         if (CallInst *CI = getCallIfRegularCall(*U, &RFI)) {
           if (!CanBeMoved(*CI))
             continue;
 
           auto Remark = [&](OptimizationRemark OR) {
-            auto newLoc = &*F.getEntryBlock().getFirstInsertionPt();
             return OR << "OpenMP runtime call "
-                      << ore::NV("OpenMPOptRuntime", RFI.Name) << " moved to "
-                      << ore::NV("OpenMPRuntimeMoves", newLoc->getDebugLoc());
+                      << ore::NV("OpenMPOptRuntime", RFI.Name)
+                      << " moved to beginning of OpenMP region";
           };
-          emitRemark<OptimizationRemark>(CI, "OpenMPRuntimeCodeMotion", Remark);
+          emitRemark<OptimizationRemark>(&F, "OpenMPRuntimeCodeMotion", Remark);
 
           CI->moveBefore(&*F.getEntryBlock().getFirstInsertionPt());
           ReplVal = CI;
           break;
         }
       if (!ReplVal)
         return false;
     }
auto ReplaceAndDeleteCB = [&](Use &U, Function &Caller) {
       if (!CI || CI == ReplVal || &F != &Caller)
         return false;
       assert(CI->getCaller() == &F && "Unexpected call!");
 
       auto Remark = [&](OptimizationRemark OR) {
         return OR << "OpenMP runtime call "
                   << ore::NV("OpenMPOptRuntime", RFI.Name) << " deduplicated";
       };
-      emitRemark<OptimizationRemark>(CI, "OpenMPRuntimeDeduplicated", Remark);
+      emitRemark<OptimizationRemark>(&F, "OpenMPRuntimeDeduplicated", Remark);
 
       CGUpdater.removeCallSite(*CI);
       CI->replaceAllUsesWith(ReplVal);
       CI->eraseFromParent();
       ++NumOpenMPRuntimeCallsDeduplicated;
       Changed = true;
       return true;
     };
private:
   /// RemarkKind should be one of the following:
   ///   - OptimizationRemark to indicate a successful optimization attempt
   ///   - OptimizationRemarkMissed to report a failed optimization attempt
   ///   - OptimizationRemarkAnalysis to provide additional information about an
   ///     optimization attempt
   ///
   /// The remark is built using a callback function provided by the caller that
   /// takes a RemarkKind as input and returns a RemarkKind.
-  template <typename RemarkKind,
-            typename RemarkCallBack = function_ref<RemarkKind(RemarkKind &&)>>
-  void emitRemark(Instruction *Inst, StringRef RemarkName,
+  template <typename RemarkKind, typename RemarkCallBack>
+  void emitRemark(Instruction *I, StringRef RemarkName,
                   RemarkCallBack &&RemarkCB) const {
-    Function *F = Inst->getParent()->getParent();
+    Function *F = I->getParent()->getParent();
     auto &ORE = OREGetter(F);
 
-    ORE.emit(
-        [&]() { return RemarkCB(RemarkKind(DEBUG_TYPE, RemarkName, Inst)); });
+    ORE.emit([&]() { return RemarkCB(RemarkKind(DEBUG_TYPE, RemarkName, I)); });
   }
 
-  /// Emit a remark on a function. Since only OptimizationRemark is supporting
-  /// this, it can't be made generic.
-  void
-  emitRemarkOnFunction(Function *F, StringRef RemarkName,
-                       function_ref<OptimizationRemark(OptimizationRemark &&)>
-                           &&RemarkCB) const {
+  /// Emit a remark on a function.
+  template <typename RemarkKind, typename RemarkCallBack>
+  void emitRemark(Function *F, StringRef RemarkName,
+                  RemarkCallBack &&RemarkCB) const {
     auto &ORE = OREGetter(F);
 
-    ORE.emit([&]() {
-      return RemarkCB(OptimizationRemark(DEBUG_TYPE, RemarkName, F));
-    });
+    ORE.emit([&]() { return RemarkCB(RemarkKind(DEBUG_TYPE, RemarkName, F)); });
   }
 
   /// The underlying module.
   Module &M;
 
   /// The SCC we are operating on.
   SmallVectorImpl<Function *> &SCC;
 
if (isKernel(F)) {
       CachedKernel = Kernel(&F);
       return *CachedKernel;
     }
 
     CachedKernel = nullptr;
     if (!F.hasLocalLinkage()) {
 
       // See https://openmp.llvm.org/remarks/OptimizationRemarks.html
-      auto Remark = [&](OptimizationRemark OR) {
-        return OR << "[OMP100] Potentially unknown OpenMP target region caller";
+      auto Remark = [&](OptimizationRemarkAnalysis ORA) {
+        return ORA
+               << "[OMP100] Potentially unknown OpenMP target region caller";
       };
-      emitRemarkOnFunction(&F, "OMP100", Remark);
+      emitRemark<OptimizationRemarkAnalysis>(&F, "OMP100", Remark);
 
       return nullptr;
     }
   }
 
   auto GetUniqueKernelForUse = [&](const Use &U) -> Kernel {
     if (auto *Cmp = dyn_cast<ICmpInst>(U.getUser())) {
       // Allow use in equality comparisons.
for (Function *F : SCC) {
     });
 
     // Do not emit a remark if we haven't seen a __kmpc_parallel_51
     // use.
     if (!KernelParallelUse)
       continue;
 
     {
-      auto Remark = [&](OptimizationRemark OR) {
-        return OR << "Found a parallel region that is called in a target "
+      auto Remark = [&](OptimizationRemarkAnalysis ORA) {
+        return ORA << "Found a parallel region that is called in a target "
                       "region but not part of a combined target construct nor "
                       "nested inside a target construct without intermediate "
                       "code. This can lead to excessive register usage for "
                       "unrelated target regions in the same translation unit "
                       "due to spurious call edges assumed by ptxas.";
       };
-      emitRemarkOnFunction(F, "OpenMPParallelRegionInNonSPMD", Remark);
+      emitRemark<OptimizationRemarkAnalysis>(F, "OpenMPParallelRegionInNonSPMD",
+                                             Remark);
     }
 
     // If this ever hits, we should investigate.
     // TODO: Checking the number of uses is not a necessary restriction and
     // should be lifted.
     if (UnknownUse || NumDirectCalls != 1 ||
         ToBeReplacedStateMachineUses.size() != 2) {
       {
-        auto Remark = [&](OptimizationRemark OR) {
-          return OR << "Parallel region is used in "
+        auto Remark = [&](OptimizationRemarkAnalysis ORA) {
+          return ORA << "Parallel region is used in "
                      << (UnknownUse ? "unknown" : "unexpected")
                      << " ways; will not attempt to rewrite the state machine.";
         };
-        emitRemarkOnFunction(F, "OpenMPParallelRegionInNonSPMD", Remark);
+        emitRemark<OptimizationRemarkAnalysis>(
+            F, "OpenMPParallelRegionInNonSPMD", Remark);
       }
       continue;
     }
 
     // Even if we have __kmpc_parallel_51 calls, we (for now) give
     // up if the function is not called from a unique kernel.
     Kernel K = getUniqueKernelFor(*F);
     if (!K) {
       {
-        auto Remark = [&](OptimizationRemark OR) {
-          return OR << "Parallel region is not known to be called from a "
+        auto Remark = [&](OptimizationRemarkAnalysis ORA) {
+          return ORA << "Parallel region is not known to be called from a "
                         "unique single target region, maybe the surrounding "
                         "function has external linkage?; will not attempt to "
                         "rewrite the state machine use.";
         };
-        emitRemarkOnFunction(F, "OpenMPParallelRegionInMultipleKernesl",
-                             Remark);
+        emitRemark<OptimizationRemarkAnalysis>(
+            F, "OpenMPParallelRegionInMultipleKernesl", Remark);
       }
       continue;
     }
 
     // We now know F is a parallel body function called only from the kernel K.
     // We also identified the state machine uses in which we replace the
     // function pointer by a new global symbol for identification purposes. This
     // ensures only direct calls to the function are left.
 
     {
-      auto RemarkParalleRegion = [&](OptimizationRemark OR) {
-        return OR << "Specialize parallel region that is only reached from a "
+      auto RemarkParalleRegion = [&](OptimizationRemarkAnalysis ORA) {
+        return ORA << "Specialize parallel region that is only reached from a "
                       "single target region to avoid spurious call edges and "
                       "excessive register usage in other target regions. "
                       "(parallel region ID: "
                    << ore::NV("OpenMPParallelRegion", F->getName())
                    << ", kernel ID: "
                    << ore::NV("OpenMPTargetRegion", K->getName()) << ")";
       };
-      emitRemarkOnFunction(F, "OpenMPParallelRegionInNonSPMD",
+      emitRemark<OptimizationRemarkAnalysis>(F, "OpenMPParallelRegionInNonSPMD",
                                              RemarkParalleRegion);
-      auto RemarkKernel = [&](OptimizationRemark OR) {
-        return OR << "Target region containing the parallel region that is "
+      auto RemarkKernel = [&](OptimizationRemarkAnalysis ORA) {
+        return ORA << "Target region containing the parallel region that is "
                       "specialized. (parallel region ID: "
                    << ore::NV("OpenMPParallelRegion", F->getName())
                    << ", kernel ID: "
                    << ore::NV("OpenMPTargetRegion", K->getName()) << ")";
       };
-      emitRemarkOnFunction(K, "OpenMPParallelRegionInNonSPMD", RemarkKernel);
+      emitRemark<OptimizationRemarkAnalysis>(K, "OpenMPParallelRegionInNonSPMD",
+                                             RemarkKernel);
     }
 
     Module &M = *F->getParent();
     Type *Int8Ty = Type::getInt8Ty(M.getContext());
 
     auto *ID = new GlobalVariable(
         M, Int8Ty, /* isConstant */ true, GlobalValue::PrivateLinkage,
         UndefValue::get(Int8Ty), F->getName() + ".ID");

llvm/test/Transforms/OpenMP/deduplication_remarks.ll

 ; 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
 @.str0 = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00", align 1
 
-; CHECK: remark: deduplication_remarks.c:9:10: OpenMP runtime call __kmpc_global_thread_num moved to deduplication_remarks.c:5:10
-; CHECK: remark: deduplication_remarks.c:7:10: OpenMP runtime call __kmpc_global_thread_num deduplicated
-; CHECK: remark: deduplication_remarks.c:5:10: OpenMP runtime call __kmpc_global_thread_num deduplicated
+; CHECK: remark: deduplication_remarks.c:4:0: OpenMP runtime call __kmpc_global_thread_num moved to beginning of OpenMP region
+; CHECK: remark: deduplication_remarks.c:4:0: OpenMP runtime call __kmpc_global_thread_num deduplicated
+; CHECK: remark: deduplication_remarks.c:4:0: OpenMP runtime call __kmpc_global_thread_num deduplicated
 define dso_local void @deduplicate() local_unnamed_addr !dbg !14 {
   %1 = tail call i32 @__kmpc_global_thread_num(%struct.ident_t* nonnull @0), !dbg !21
   call void @useI32(i32 %1), !dbg !23
   %2 = tail call i32 @__kmpc_global_thread_num(%struct.ident_t* nonnull @0), !dbg !24
   call void @useI32(i32 %2), !dbg !25
   %3 = tail call i32 @__kmpc_global_thread_num(%struct.ident_t* nonnull @0), !dbg !26
   call void @useI32(i32 %3), !dbg !27
   ret void, !dbg !28

llvm/test/Transforms/OpenMP/gpu_kernel_detection_remarks.ll

-; 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
+; RUN: opt -passes=openmp-opt-cgscc -pass-remarks-analysis=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-analysis=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/icv_remarks.ll

-; 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
+; RUN: opt -passes=openmp-opt-cgscc -pass-remarks-analysis=openmp-opt -openmp-print-icv-values -disable-output < %s 2>&1 | FileCheck %s
+; RUN: opt -openmp-opt-cgscc -pass-remarks-analysis=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* }