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 @@ -1587,7 +1587,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. @@ -1599,7 +1605,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 @@ -949,8 +949,7 @@ return OR << "."; }; - emitRemark(MergableCIs.front(), - "OpenMPParallelRegionMerging", Remark); + emitRemark(MergableCIs.front(), "OMP150", Remark); Function *OriginalFn = BB->getParent(); LLVM_DEBUG(dbgs() << TAG << "Merge " << MergableCIs.size() @@ -1204,8 +1203,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(); @@ -1311,7 +1309,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; @@ -1593,9 +1591,9 @@ << ore::NV("OpenMPOptRuntime", RFI.Name) << " deduplicated."; }; if (CI->getDebugLoc()) - emitRemark(CI, "OpenMPRuntimeDeduplicated", Remark); + emitRemark(CI, "OMP170", Remark); else - emitRemark(&F, "OpenMPRuntimeDeduplicated", Remark); + emitRemark(&F, "OMP170", Remark); CGUpdater.removeCallSite(*CI); CI->replaceAllUsesWith(ReplVal); @@ -1702,7 +1700,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. @@ -1711,7 +1716,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. @@ -1880,8 +1892,7 @@ << (UnknownUse ? "unknown" : "unexpected") << " ways. Will not attempt to rewrite the state machine."; }; - emitRemark(F, "OpenMPParallelRegionInNonSPMD", - Remark); + emitRemark(F, "OMP101", Remark); continue; } @@ -1893,8 +1904,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; } @@ -2581,8 +2591,7 @@ << ((AllocSize->getZExtValue() != 1) ? " bytes " : " byte ") << "of shared memory."; }; - A.emitRemark(CB, "OpenMPReplaceGlobalization", - Remark); + A.emitRemark(CB, "OMP111", Remark); SharedMem->setAlignment(MaybeAlign(32)); @@ -2823,8 +2832,8 @@ } return ORA << "."; }; - A.emitRemark( - NonCompatibleI, "OpenMPKernelNonSPMDMode", Remark); + A.emitRemark(NonCompatibleI, "OMP121", + Remark); LLVM_DEBUG(dbgs() << TAG << "SPMD-incompatible side-effect: " << *NonCompatibleI << "\n"); @@ -2864,8 +2873,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; }; @@ -2909,8 +2917,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; } @@ -2923,8 +2930,7 @@ return OR << "Rewriting generic-mode kernel with a customized state " "machine."; }; - A.emitRemark( - KernelInitCB, "OpenMPKernelWithCustomizedStateMachine", Remark); + A.emitRemark(KernelInitCB, "OMP131", Remark); } else { ++NumOpenMPTargetRegionKernelsCustomStateMachineWithFallback; @@ -2932,9 +2938,7 @@ return OR << "Generic-mode kernel is executed with a customized state " "machine that requires a fallback."; }; - A.emitRemark( - KernelInitCB, "OpenMPKernelWithCustomizedStateMachineAndFallback", - Remark); + A.emitRemark(KernelInitCB, "OMP132", Remark); // Tell the user why we ended up with a fallback. for (CallBase *UnknownParallelRegionCB : ReachedUnknownParallelRegions) { @@ -2945,9 +2949,8 @@ << "`__attribute__((assume(\"omp_no_parallelism\")))` to " "override."; }; - A.emitRemark( - UnknownParallelRegionCB, - "OpenMPKernelWithCustomizedStateMachineAndFallback", Remark); + A.emitRemark(UnknownParallelRegionCB, + "OMP133", Remark); } } @@ -3741,7 +3744,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."; });