diff --git a/clang/test/OpenMP/remarks_parallel_in_multiple_target_state_machines.c b/clang/test/OpenMP/remarks_parallel_in_multiple_target_state_machines.c --- a/clang/test/OpenMP/remarks_parallel_in_multiple_target_state_machines.c +++ b/clang/test/OpenMP/remarks_parallel_in_multiple_target_state_machines.c @@ -8,23 +8,23 @@ 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 { } @@ -37,9 +37,9 @@ 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 { } @@ -55,9 +55,9 @@ 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 { } @@ -83,4 +83,4 @@ } } -// all-remark@* 9 {{OpenMP runtime call __kmpc_global_thread_num deduplicated}} +// all-remark@* 9 {{OpenMP runtime call __kmpc_global_thread_num deduplicated. [OMP170]}} diff --git a/clang/test/OpenMP/remarks_parallel_in_target_state_machine.c b/clang/test/OpenMP/remarks_parallel_in_target_state_machine.c --- a/clang/test/OpenMP/remarks_parallel_in_target_state_machine.c +++ b/clang/test/OpenMP/remarks_parallel_in_target_state_machine.c @@ -8,16 +8,16 @@ 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 { } @@ -40,4 +40,4 @@ } } -// expected-remark@* {{OpenMP runtime call __kmpc_global_thread_num deduplicated}} +// expected-remark@* {{OpenMP runtime call __kmpc_global_thread_num deduplicated. [OMP170]}} diff --git a/llvm/include/llvm/Transforms/IPO/Attributor.h b/llvm/include/llvm/Transforms/IPO/Attributor.h --- a/llvm/include/llvm/Transforms/IPO/Attributor.h +++ b/llvm/include/llvm/Transforms/IPO/Attributor.h @@ -1585,7 +1585,13 @@ Function *F = I->getFunction(); auto &ORE = OREGetter.getValue()(F); - ORE.emit([&]() { return RemarkCB(RemarkKind(PassName, RemarkName, I)); }); + 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. @@ -1597,7 +1603,13 @@ auto &ORE = OREGetter.getValue()(F); - ORE.emit([&]() { return RemarkCB(RemarkKind(PassName, RemarkName, 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) diff --git a/llvm/lib/Transforms/IPO/AttributorAttributes.cpp b/llvm/lib/Transforms/IPO/AttributorAttributes.cpp --- a/llvm/lib/Transforms/IPO/AttributorAttributes.cpp +++ b/llvm/lib/Transforms/IPO/AttributorAttributes.cpp @@ -5039,7 +5039,10 @@ return OR << "Moving globalized variable to the stack."; return OR << "Moving memory allocation from the heap to the stack."; }; - A.emitRemark(AI.CB, "HeapToStack", Remark); + if (AI.LibraryFunctionId == LibFunc___kmpc_alloc_shared) + A.emitRemark(AI.CB, "OMP110", Remark); + else + A.emitRemark(AI.CB, "HeapToStack", Remark); Value *Size; Optional SizeAPI = getSize(A, *this, AI); @@ -5335,8 +5338,7 @@ }; if (AI.LibraryFunctionId == LibFunc___kmpc_alloc_shared) - A.emitRemark(AI.CB, "HeapToStackFailed", - Remark); + A.emitRemark(AI.CB, "OMP113", Remark); LLVM_DEBUG(dbgs() << "[H2S] Bad user: " << *UserI << "\n"); ValidUsesOnly = false; diff --git a/llvm/lib/Transforms/IPO/OpenMPOpt.cpp b/llvm/lib/Transforms/IPO/OpenMPOpt.cpp --- a/llvm/lib/Transforms/IPO/OpenMPOpt.cpp +++ b/llvm/lib/Transforms/IPO/OpenMPOpt.cpp @@ -943,8 +943,7 @@ return OR << "."; }; - emitRemark(MergableCIs.front(), - "OpenMPParallelRegionMerging", Remark); + emitRemark(MergableCIs.front(), "OMP150", Remark); Function *OriginalFn = BB->getParent(); LLVM_DEBUG(dbgs() << TAG << "Merge " << MergableCIs.size() @@ -1040,8 +1039,7 @@ << "."; }; if (CI != MergableCIs.front()) - emitRemark(CI, "OpenMPParallelRegionMerging", - Remark); + emitRemark(CI, "OMP150", Remark); CI->eraseFromParent(); } @@ -1210,8 +1208,7 @@ auto Remark = [&](OptimizationRemark OR) { return OR << "Removing parallel region with no side-effects."; }; - emitRemark(CI, "OpenMPParallelRegionDeletion", - Remark); + emitRemark(CI, "OMP160", Remark); CGUpdater.removeCallSite(*CI); CI->eraseFromParent(); @@ -1317,7 +1314,7 @@ << "Found thread data sharing on the GPU. " << "Expect degraded performance due to data globalization."; }; - emitRemark(CI, "OpenMPGlobalization", Remark); + emitRemark(CI, "OMP112", Remark); } return false; @@ -1598,7 +1595,7 @@ return OR << "OpenMP runtime call " << ore::NV("OpenMPOptRuntime", RFI.Name) << " deduplicated."; }; - emitRemark(&F, "OpenMPRuntimeDeduplicated", Remark); + emitRemark(&F, "OMP170", Remark); CGUpdater.removeCallSite(*CI); CI->replaceAllUsesWith(ReplVal); @@ -1705,7 +1702,14 @@ 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. @@ -1714,7 +1718,14 @@ 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. @@ -1883,8 +1894,7 @@ << (UnknownUse ? "unknown" : "unexpected") << " ways. Will not attempt to rewrite the state machine."; }; - emitRemark(F, "OpenMPParallelRegionInNonSPMD", - Remark); + emitRemark(F, "OMP101", Remark); continue; } @@ -1896,8 +1906,7 @@ return ORA << "Parallel region is not called from a unique kernel. " "Will not attempt to rewrite the state machine."; }; - emitRemark( - F, "OpenMPParallelRegionInMultipleKernesl", Remark); + emitRemark(F, "OMP102", Remark); continue; } @@ -2584,8 +2593,7 @@ << ((AllocSize->getZExtValue() != 1) ? " bytes " : " byte ") << "of shared memory"; }; - A.emitRemark(CB, "OpenMPReplaceGlobalization", - Remark); + A.emitRemark(CB, "OMP111", Remark); SharedMem->setAlignment(MaybeAlign(32)); @@ -2816,14 +2824,14 @@ auto Remark = [&](OptimizationRemarkAnalysis ORA) { ORA << "Value has potential side effects preventing SPMD-mode " "execution"; - if (auto *CI = dyn_cast(NonCompatibleI)) { + if (isa(NonCompatibleI)) { ORA << ". Add `__attribute__((assume(\"ompx_spmd_amenable\")))` to " "the called function to override"; } return ORA << "."; }; - A.emitRemark( - NonCompatibleI, "OpenMPKernelNonSPMDMode", Remark); + A.emitRemark(NonCompatibleI, "OMP121", + Remark); LLVM_DEBUG(dbgs() << TAG << "SPMD-incompatible side-effect: " << *NonCompatibleI << "\n"); @@ -2863,8 +2871,7 @@ auto Remark = [&](OptimizationRemark OR) { return OR << "Transformed Generic-mode kernel to SPMD-mode."; }; - A.emitRemark(KernelInitCB, "OpenMPKernelSPMDMode", - Remark); + A.emitRemark(KernelInitCB, "OMP120", Remark); return true; }; @@ -2908,8 +2915,7 @@ auto Remark = [&](OptimizationRemark OR) { return OR << "Removing unused state machine from generic-mode kernel."; }; - A.emitRemark( - KernelInitCB, "OpenMPKernelWithoutStateMachine", Remark); + A.emitRemark(KernelInitCB, "OMP130", Remark); return ChangeStatus::CHANGED; } @@ -2922,8 +2928,7 @@ return OR << "Rewriting generic-mode kernel with a customized state " "machine."; }; - A.emitRemark( - KernelInitCB, "OpenMPKernelWithCustomizedStateMachine", Remark); + A.emitRemark(KernelInitCB, "OMP131", Remark); } else { ++NumOpenMPTargetRegionKernelsCustomStateMachineWithFallback; @@ -2937,9 +2942,7 @@ ReachedUnknownParallelRegions.size()) << " unkown parallel regions]."; }; - A.emitRemark( - KernelInitCB, "OpenMPKernelWithCustomizedStateMachineAndFallback", - Remark); + A.emitRemark(KernelInitCB, "OMP132", Remark); // Tell the user why we ended up with a fallback. for (CallBase *UnknownParallelRegionCB : ReachedUnknownParallelRegions) { @@ -2950,9 +2953,8 @@ << "`__attribute__((assume(\"omp_no_parallelism\")))` to " "override."; }; - A.emitRemark( - UnknownParallelRegionCB, - "OpenMPKernelWithCustomizedStateMachineAndFallback", Remark); + A.emitRemark(UnknownParallelRegionCB, + "OMP133", Remark); } } @@ -3501,7 +3503,7 @@ auto EmitRemark = [&](Function &F) { auto &ORE = FAM.getResult(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."; });