[OpenMP] Add IDs to OpenMP remarks

[jhuber6]

Summary

This patch adds unique idenfitiers to the existing OpenMP remarks. This makes
it easier to identify the corresponding documentation for each remark that will
be hosted in the OpenMP webpage.

Depends on D105898

[jdoerfert]

Don't any tests change?

[jhuber6]

In D105939#2875858, @jdoerfert wrote:

Don't any tests change?

The reamrk ID is postfixed, so the regexes should still match.

[jhuber6]

Adding test case for IDs.

[jdoerfert]

LG

clang/test/OpenMP/remarks_parallel_in_multiple_target_state_machines.c

 // RUN: %clang_cc1                                 -verify=host      -Rpass=openmp-opt -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=all,safe  -Rpass=openmp-opt -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=all,safe  -Rpass=openmp-opt -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 baz(void) __attribute__((assume("omp_no_openmp")));
 
 void bar1(void) {
 #pragma omp parallel // #0
-                     // safe-remark@#0 {{Parallel region is used in unknown ways. Will not attempt to rewrite the state machine.}}
+                     // safe-remark@#0 {{Parallel region is used in unknown ways. Will not attempt to rewrite the state machine. [OMP101]}}
   {
   }
 }
 void bar2(void) {
 #pragma omp parallel // #1
-                     // safe-remark@#1 {{Parallel region is used in unknown ways. Will not attempt to rewrite the state machine.}}
+                     // safe-remark@#1 {{Parallel region is used in unknown ways. Will not attempt to rewrite the state machine. [OMP101]}}
   {
   }
 }
 
 void foo1(void) {
 #pragma omp target teams // #2
-                         // all-remark@#2 {{Rewriting generic-mode kernel with a customized state machine.}}
+                         // all-remark@#2 {{Rewriting generic-mode kernel with a customized state machine. [OMP131]}}
 
   {
-    baz();           // all-remark {{Value has potential side effects preventing SPMD-mode execution. Add `__attribute__((assume("ompx_spmd_amenable")))` to the called function to override.}}
+    baz();           // all-remark {{Value has potential side effects preventing SPMD-mode execution. Add `__attribute__((assume("ompx_spmd_amenable")))` to the called function to override. [OMP121]}}
 #pragma omp parallel // #3
     {
     }
     bar1();
 #pragma omp parallel // #4
     {
     }
   }
 }
 
 void foo2(void) {
 #pragma omp target teams // #5
-                         // all-remark@#5 {{Rewriting generic-mode kernel with a customized state machine.}}
+                         // all-remark@#5 {{Rewriting generic-mode kernel with a customized state machine. [OMP131]}}
   {
-    baz();           // all-remark {{Value has potential side effects preventing SPMD-mode execution. Add `__attribute__((assume("ompx_spmd_amenable")))` to the called function to override.}}
+    baz();           // all-remark {{Value has potential side effects preventing SPMD-mode execution. Add `__attribute__((assume("ompx_spmd_amenable")))` to the called function to override. [OMP121]}}
 #pragma omp parallel // #6
     {
     }
     bar1();
     bar2();
 #pragma omp parallel // #7
     {
     }
     bar1();
     bar2();
   }
 }
 
 void foo3(void) {
 #pragma omp target teams // #8
-                         // all-remark@#8 {{Rewriting generic-mode kernel with a customized state machine.}}
+                         // all-remark@#8 {{Rewriting generic-mode kernel with a customized state machine. [OMP131]}}
   {
-    baz();           // all-remark {{Value has potential side effects preventing SPMD-mode execution. Add `__attribute__((assume("ompx_spmd_amenable")))` to the called function to override.}}
+    baz();           // all-remark {{Value has potential side effects preventing SPMD-mode execution. Add `__attribute__((assume("ompx_spmd_amenable")))` to the called function to override. [OMP121]}}
 #pragma omp parallel // #9
     {
     }
     bar1();
     bar2();
 #pragma omp parallel // #10
     {
     }
#pragma omp parallel
   {
   }
 
 #pragma omp target teams distribute parallel for
   for (int i = 0; i < 100; ++i) {
   }
 }
 
-// all-remark@* 9 {{OpenMP runtime call __kmpc_global_thread_num deduplicated}}
+// all-remark@* 9 {{OpenMP runtime call __kmpc_global_thread_num deduplicated. [OMP170]}}

clang/test/OpenMP/remarks_parallel_in_target_state_machine.c

 // RUN: %clang_cc1                                 -verify=host -Rpass=openmp-opt -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-opt -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-opt -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 baz(void) __attribute__((assume("omp_no_openmp")));
 
 void bar(void) {
 #pragma omp parallel // #1                                                                                                                                                                                                                                                                                                                                           \
-                     // expected-remark@#1 {{Parallel region is used in unknown ways. Will not attempt to rewrite the state machine.}}
+                     // expected-remark@#1 {{Parallel region is used in unknown ways. Will not attempt to rewrite the state machine. [OMP101]}}
   {
   }
 }
 
 void foo(void) {
 #pragma omp target teams // #2
-                         // expected-remark@#2 {{Rewriting generic-mode kernel with a customized state machine.}}
+                         // expected-remark@#2 {{Rewriting generic-mode kernel with a customized state machine. [OMP131]}}
   {
-    baz();               // expected-remark {{Value has potential side effects preventing SPMD-mode execution. Add `__attribute__((assume("ompx_spmd_amenable")))` to the called function to override.}}
+    baz();               // expected-remark {{Value has potential side effects preventing SPMD-mode execution. Add `__attribute__((assume("ompx_spmd_amenable")))` to the called function to override. [OMP121]}}
 #pragma omp parallel
     {
     }
     bar();
 #pragma omp parallel
     {
     }
   }
 }
 
 void spmd(void) {
   // Verify we do not emit the remarks above for "SPMD" regions.
 #pragma omp target teams
 #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 deduplicated}}
+// expected-remark@* {{OpenMP runtime call __kmpc_global_thread_num deduplicated. [OMP170]}}

llvm/include/llvm/Transforms/IPO/Attributor.h

public:
   void emitRemark(Instruction *I, StringRef RemarkName,
                   RemarkCallBack &&RemarkCB) const {
     if (!OREGetter)
       return;
 
     Function *F = I->getFunction();
     auto &ORE = OREGetter.getValue()(F);
 
+    if (RemarkName.startswith("OMP"))
+      ORE.emit([&]() {
+        return RemarkCB(RemarkKind(PassName, RemarkName, I))
+               << " [" << RemarkName << "]";
+      });
+    else
       ORE.emit([&]() { return RemarkCB(RemarkKind(PassName, RemarkName, I)); });
   }
 
   /// Emit a remark on a function.
   template <typename RemarkKind, typename RemarkCallBack>
   void emitRemark(Function *F, StringRef RemarkName,
                   RemarkCallBack &&RemarkCB) const {
     if (!OREGetter)
       return;
 
     auto &ORE = OREGetter.getValue()(F);
 
+    if (RemarkName.startswith("OMP"))
+      ORE.emit([&]() {
+        return RemarkCB(RemarkKind(PassName, RemarkName, F))
+               << " [" << RemarkName << "]";
+      });
+    else
       ORE.emit([&]() { return RemarkCB(RemarkKind(PassName, RemarkName, F)); });
   }
 
   /// Helper struct used in the communication between an abstract attribute (AA)
   /// that wants to change the signature of a function and the Attributor which
   /// applies the changes. The struct is partially initialized with the
   /// information from the AA (see the constructor). All other members are
   /// provided by the Attributor prior to invoking any callbacks.
   struct ArgumentReplacementInfo {

llvm/lib/Transforms/IPO/AttributorAttributes.cpp

for (auto &It : AllocationInfos) {
 
       auto Remark = [&](OptimizationRemark OR) {
         LibFunc IsAllocShared;
         if (TLI->getLibFunc(*AI.CB, IsAllocShared))
           if (IsAllocShared == LibFunc___kmpc_alloc_shared)
             return OR << "Moving globalized variable to the stack.";
         return OR << "Moving memory allocation from the heap to the stack.";
       };
+      if (AI.LibraryFunctionId == LibFunc___kmpc_alloc_shared)
+        A.emitRemark<OptimizationRemark>(AI.CB, "OMP110", Remark);
+      else
         A.emitRemark<OptimizationRemark>(AI.CB, "HeapToStack", Remark);
 
       Value *Size;
       Optional<APInt> SizeAPI = getSize(A, *this, AI);
       if (SizeAPI.hasValue()) {
         Size = ConstantInt::get(AI.CB->getContext(), *SizeAPI);
       } else if (AI.Kind == AllocationInfo::AllocationKind::CALLOC) {
         auto *Num = AI.CB->getOperand(0);
         auto *SizeT = AI.CB->getOperand(1);
auto Pred = [&](const Use &U, bool &Follow) -> bool {
           auto Remark = [&](OptimizationRemarkMissed ORM) {
             return ORM
                    << "Could not move globalized variable to the stack. "
                       "Variable is potentially captured in call. Mark "
                       "parameter as `__attribute__((noescape))` to override.";
           };
 
           if (AI.LibraryFunctionId == LibFunc___kmpc_alloc_shared)
-            A.emitRemark<OptimizationRemarkMissed>(AI.CB, "HeapToStackFailed",
-                                                   Remark);
+            A.emitRemark<OptimizationRemarkMissed>(AI.CB, "OMP113", Remark);
 
           LLVM_DEBUG(dbgs() << "[H2S] Bad user: " << *UserI << "\n");
           ValidUsesOnly = false;
         }
         return true;
       }
 
       if (isa<GetElementPtrInst>(UserI) || isa<BitCastInst>(UserI) ||

llvm/lib/Transforms/IPO/OpenMPOpt.cpp

auto Merge = [&](SmallVectorImpl<CallInst *> &MergableCIs, BasicBlock *BB) {
         for (auto *CI : llvm::drop_begin(MergableCIs)) {
           OR << ore::NV("OpenMPParallelMerge", CI->getDebugLoc());
           if (CI != MergableCIs.back())
             OR << ", ";
         }
         return OR << ".";
       };
 
-      emitRemark<OptimizationRemark>(MergableCIs.front(),
-                                     "OpenMPParallelRegionMerging", Remark);
+      emitRemark<OptimizationRemark>(MergableCIs.front(), "OMP150", Remark);
 
       Function *OriginalFn = BB->getParent();
       LLVM_DEBUG(dbgs() << TAG << "Merge " << MergableCIs.size()
                         << " parallel regions in " << OriginalFn->getName()
                         << "\n");
 
       // Isolate the calls to merge in a separate block.
       EndBB = SplitBlock(BB, MergableCIs.back()->getNextNode(), DT, LI);
auto Merge = [&](SmallVectorImpl<CallInst *> &MergableCIs, BasicBlock *BB) {
           return OR << "Parallel region at "
                     << ore::NV("OpenMPParallelMerge", CI->getDebugLoc())
                     << " merged with "
                     << ore::NV("OpenMPParallelMergeFront",
                                MergableCIs.front()->getDebugLoc())
                     << ".";
         };
         if (CI != MergableCIs.front())
-          emitRemark<OptimizationRemark>(CI, "OpenMPParallelRegionMerging",
-                                         Remark);
+          emitRemark<OptimizationRemark>(CI, "OMP150", Remark);
 
         CI->eraseFromParent();
       }
 
       assert(OutlinedFn != OriginalFn && "Outlining failed");
       CGUpdater.registerOutlinedFunction(*OriginalFn, *OutlinedFn);
       CGUpdater.reanalyzeFunction(*OriginalFn);
 
auto DeleteCallCB = [&](Use &U, Function &) {
         return false;
 
       LLVM_DEBUG(dbgs() << TAG << "Delete read-only parallel region in "
                         << CI->getCaller()->getName() << "\n");
 
       auto Remark = [&](OptimizationRemark OR) {
         return OR << "Removing parallel region with no side-effects.";
       };
-      emitRemark<OptimizationRemark>(CI, "OpenMPParallelRegionDeletion",
-                                     Remark);
+      emitRemark<OptimizationRemark>(CI, "OMP160", Remark);
 
       CGUpdater.removeCallSite(*CI);
       CI->eraseFromParent();
       Changed = true;
       ++NumOpenMPParallelRegionsDeleted;
       return true;
     };
 
void analysisGlobalization() {
 
     auto CheckGlobalization = [&](Use &U, Function &Decl) {
       if (CallInst *CI = getCallIfRegularCall(U, &RFI)) {
         auto Remark = [&](OptimizationRemarkMissed ORM) {
           return ORM
                  << "Found thread data sharing on the GPU. "
                  << "Expect degraded performance due to data globalization.";
         };
-        emitRemark<OptimizationRemarkMissed>(CI, "OpenMPGlobalization", Remark);
+        emitRemark<OptimizationRemarkMissed>(CI, "OMP112", Remark);
       }
 
       return false;
     };
 
     RFI.foreachUse(SCC, CheckGlobalization);
   }
 
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>(&F, "OpenMPRuntimeDeduplicated", Remark);
+      emitRemark<OptimizationRemark>(&F, "OMP170", Remark);
 
       CGUpdater.removeCallSite(*CI);
       CI->replaceAllUsesWith(ReplVal);
       CI->eraseFromParent();
       ++NumOpenMPRuntimeCallsDeduplicated;
       Changed = true;
       return true;
     };
private:
   /// 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>
   void emitRemark(Instruction *I, StringRef RemarkName,
                   RemarkCallBack &&RemarkCB) const {
     Function *F = I->getParent()->getParent();
     auto &ORE = OREGetter(F);
 
-    ORE.emit([&]() { return RemarkCB(RemarkKind(DEBUG_TYPE, RemarkName, I)); });
+    if (RemarkName.startswith("OMP"))
+      ORE.emit([&]() {
+        return RemarkCB(RemarkKind(DEBUG_TYPE, RemarkName, I))
+               << " [" << RemarkName << "]";
+      });
+    else
+      ORE.emit(
+          [&]() { return RemarkCB(RemarkKind(DEBUG_TYPE, RemarkName, I)); });
   }
 
   /// 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(RemarkKind(DEBUG_TYPE, RemarkName, F)); });
+    if (RemarkName.startswith("OMP"))
+      ORE.emit([&]() {
+        return RemarkCB(RemarkKind(DEBUG_TYPE, RemarkName, F))
+               << " [" << RemarkName << "]";
+      });
+    else
+      ORE.emit(
+          [&]() { return RemarkCB(RemarkKind(DEBUG_TYPE, RemarkName, F)); });
   }
 
   /// The underlying module.
   Module &M;
 
   /// The SCC we are operating on.
   SmallVectorImpl<Function *> &SCC;
 
for (Function *F : SCC) {
     // should be lifted.
     if (UnknownUse || NumDirectCalls != 1 ||
         ToBeReplacedStateMachineUses.size() > 2) {
       auto Remark = [&](OptimizationRemarkAnalysis ORA) {
         return ORA << "Parallel region is used in "
                    << (UnknownUse ? "unknown" : "unexpected")
                    << " ways. Will not attempt to rewrite the state machine.";
       };
-      emitRemark<OptimizationRemarkAnalysis>(F, "OpenMPParallelRegionInNonSPMD",
-                                             Remark);
+      emitRemark<OptimizationRemarkAnalysis>(F, "OMP101", 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 = [&](OptimizationRemarkAnalysis ORA) {
         return ORA << "Parallel region is not called from a unique kernel. "
                       "Will not attempt to rewrite the state machine.";
       };
-      emitRemark<OptimizationRemarkAnalysis>(
-          F, "OpenMPParallelRegionInMultipleKernesl", Remark);
+      emitRemark<OptimizationRemarkAnalysis>(F, "OMP102", 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.
 
for (CallBase *CB : MallocCalls) {
           ConstantExpr::getPointerCast(SharedMem, Int8Ty->getPointerTo());
 
       auto Remark = [&](OptimizationRemark OR) {
         return OR << "Replaced globalized variable with "
                   << ore::NV("SharedMemory", AllocSize->getZExtValue())
                   << ((AllocSize->getZExtValue() != 1) ? " bytes " : " byte ")
                   << "of shared memory";
       };
-      A.emitRemark<OptimizationRemark>(CB, "OpenMPReplaceGlobalization",
-                                       Remark);
+      A.emitRemark<OptimizationRemark>(CB, "OMP111", Remark);
 
       SharedMem->setAlignment(MaybeAlign(32));
 
       A.changeValueAfterManifest(*CB, *NewBuffer);
       A.deleteAfterManifest(*CB);
       A.deleteAfterManifest(*FreeCalls.front());
 
       NumBytesMovedToSharedMemory += AllocSize->getZExtValue();
if (!SPMDCompatibilityTracker.isAssumed()) {
         // Skip diagnostics on calls to known OpenMP runtime functions for now.
         if (auto *CB = dyn_cast<CallBase>(NonCompatibleI))
           if (OMPInfoCache.RTLFunctions.contains(CB->getCalledFunction()))
             continue;
 
         auto Remark = [&](OptimizationRemarkAnalysis ORA) {
           ORA << "Value has potential side effects preventing SPMD-mode "
                  "execution";
-          if (auto *CI = dyn_cast<CallBase>(NonCompatibleI)) {
+          if (isa<CallBase>(NonCompatibleI)) {
             ORA << ". Add `__attribute__((assume(\"ompx_spmd_amenable\")))` to "
                    "the called function to override";
           }
           return ORA << ".";
         };
-        A.emitRemark<OptimizationRemarkAnalysis>(
-            NonCompatibleI, "OpenMPKernelNonSPMDMode", Remark);
+        A.emitRemark<OptimizationRemarkAnalysis>(NonCompatibleI, "OMP121",
+                                                 Remark);
 
         LLVM_DEBUG(dbgs() << TAG << "SPMD-incompatible side-effect: "
                           << *NonCompatibleI << "\n");
       }
 
       return false;
     }
 
bool changeToSPMDMode(Attributor &A) {
     A.changeUseAfterManifest(
         KernelDeinitCB->getArgOperandUse(DeinitIsSPMDArgNo),
         *ConstantInt::getBool(Ctx, 1));
     ++NumOpenMPTargetRegionKernelsSPMD;
 
     auto Remark = [&](OptimizationRemark OR) {
       return OR << "Transformed Generic-mode kernel to SPMD-mode.";
     };
-    A.emitRemark<OptimizationRemark>(KernelInitCB, "OpenMPKernelSPMDMode",
-                                     Remark);
+    A.emitRemark<OptimizationRemark>(KernelInitCB, "OMP120", Remark);
     return true;
   };
 
   ChangeStatus buildCustomStateMachine(Attributor &A) {
     assert(ReachedKnownParallelRegions.isValidState() &&
            "Custom state machine with invalid parallel region states?");
 
     const int InitIsSPMDArgNo = 1;
ChangeStatus buildCustomStateMachine(Attributor &A) {
     // to do the work alone.
     if (ReachedKnownParallelRegions.empty() &&
         ReachedUnknownParallelRegions.empty()) {
       ++NumOpenMPTargetRegionKernelsWithoutStateMachine;
 
       auto Remark = [&](OptimizationRemark OR) {
         return OR << "Removing unused state machine from generic-mode kernel.";
       };
-      A.emitRemark<OptimizationRemark>(
-          KernelInitCB, "OpenMPKernelWithoutStateMachine", Remark);
+      A.emitRemark<OptimizationRemark>(KernelInitCB, "OMP130", Remark);
 
       return ChangeStatus::CHANGED;
     }
 
     // Keep track in the statistics of our new shiny custom state machine.
     if (ReachedUnknownParallelRegions.empty()) {
       ++NumOpenMPTargetRegionKernelsCustomStateMachineWithoutFallback;
 
       auto Remark = [&](OptimizationRemark OR) {
         return OR << "Rewriting generic-mode kernel with a customized state "
                      "machine.";
       };
-      A.emitRemark<OptimizationRemark>(
-          KernelInitCB, "OpenMPKernelWithCustomizedStateMachine", Remark);
+      A.emitRemark<OptimizationRemark>(KernelInitCB, "OMP131", Remark);
     } else {
       ++NumOpenMPTargetRegionKernelsCustomStateMachineWithFallback;
 
       auto Remark = [&](OptimizationRemarkAnalysis OR) {
         return OR << "Generic-mode kernel is executed with a customized state "
                      "machine that requires a fallback ["
                   << ore::NV("ParallelRegions",
                              ReachedKnownParallelRegions.size())
                   << " known parallel regions, "
                   << ore::NV("UnknownParallelRegions",
                              ReachedUnknownParallelRegions.size())
                   << " unkown parallel regions].";
       };
-      A.emitRemark<OptimizationRemarkAnalysis>(
-          KernelInitCB, "OpenMPKernelWithCustomizedStateMachineAndFallback",
-          Remark);
+      A.emitRemark<OptimizationRemarkAnalysis>(KernelInitCB, "OMP132", Remark);
 
       // Tell the user why we ended up with a fallback.
       for (CallBase *UnknownParallelRegionCB : ReachedUnknownParallelRegions) {
         if (!UnknownParallelRegionCB)
           continue;
         auto Remark = [&](OptimizationRemarkAnalysis ORA) {
           return ORA << "Call may contain unknown parallel regions. Use "
                      << "`__attribute__((assume(\"omp_no_parallelism\")))` to "
                         "override.";
         };
-        A.emitRemark<OptimizationRemarkAnalysis>(
-            UnknownParallelRegionCB,
-            "OpenMPKernelWithCustomizedStateMachineAndFallback", Remark);
+        A.emitRemark<OptimizationRemarkAnalysis>(UnknownParallelRegionCB,
+                                                 "OMP133", Remark);
       }
     }
 
     // Create all the blocks:
     //
     //                       InitCB = __kmpc_target_init(...)
     //                       bool IsWorker = InitCB >= 0;
     //                       if (IsWorker) {
for (const User *U : F.users())
       if (!isa<BlockAddress>(U))
         return true;
     return false;
   };
 
   auto EmitRemark = [&](Function &F) {
     auto &ORE = FAM.getResult<OptimizationRemarkEmitterAnalysis>(F);
     ORE.emit([&]() {
-      OptimizationRemarkAnalysis ORA(DEBUG_TYPE, "InternalizationFailure", &F);
+      OptimizationRemarkAnalysis ORA(DEBUG_TYPE, "OMP140", &F);
       return ORA << "Could not internalize function. "
                  << "Some optimizations may not be possible.";
     });
   };
 
   // Create internal copies of each function if this is a kernel Module. This
   // allows iterprocedural passes to see every call edge.
   DenseSet<const Function *> InternalizedFuncs;