[Offload][OpenMP] Prettify error messages by "demangling" the kernel name #101400
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
|
@llvm/pr-subscribers-offload @llvm/pr-subscribers-flang-openmp Author: Johannes Doerfert (jdoerfert) ChangesThe kernel names for OpenMP are manually mangled and not ideal when we report something to the user. We demangle them now, providing the function and line number of the target region, together with the actual kernel name. Full diff: https://github.com/llvm/llvm-project/pull/101400.diff 14 Files Affected:
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMP.h b/llvm/include/llvm/Frontend/OpenMP/OMP.h
index 6f7a39acac1d3..f081015db0b0b 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMP.h
+++ b/llvm/include/llvm/Frontend/OpenMP/OMP.h
@@ -17,6 +17,7 @@
#include "llvm/ADT/ArrayRef.h"
#include "llvm/ADT/SmallVector.h"
+#include "llvm/ADT/StringRef.h"
namespace llvm::omp {
ArrayRef<Directive> getLeafConstructs(Directive D);
@@ -30,6 +31,14 @@ Directive getCompoundConstruct(ArrayRef<Directive> Parts);
bool isLeafConstruct(Directive D);
bool isCompositeConstruct(Directive D);
bool isCombinedConstruct(Directive D);
+
+/// Create a nicer version of a function name for humans to look at.
+std::string prettityFunctionName(StringRef FunctionName);
+
+/// Deconstruct an OpenMP kernel name into the parent function name and the line
+/// number.
+std::string deconstructOpenMPKernelName(StringRef KernelName, unsigned &LineNo);
+
} // namespace llvm::omp
#endif // LLVM_FRONTEND_OPENMP_OMP_H
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
index 1614d5716d28c..9cb311834907b 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
@@ -196,6 +196,9 @@ class OpenMPIRBuilderConfig {
/// Data structure to contain the information needed to uniquely identify
/// a target entry.
struct TargetRegionEntryInfo {
+ /// The prefix used for kernel names.
+ static constexpr const char *KernelNamePrefix = "__omp_offloading_";
+
std::string ParentName;
unsigned DeviceID;
unsigned FileID;
diff --git a/llvm/lib/Frontend/OpenMP/CMakeLists.txt b/llvm/lib/Frontend/OpenMP/CMakeLists.txt
index 67aedf5c2b61a..82d2a9ae7c533 100644
--- a/llvm/lib/Frontend/OpenMP/CMakeLists.txt
+++ b/llvm/lib/Frontend/OpenMP/CMakeLists.txt
@@ -17,6 +17,7 @@ add_llvm_component_library(LLVMFrontendOpenMP
TargetParser
TransformUtils
Analysis
+ Demangle
MC
Scalar
BitReader
diff --git a/llvm/lib/Frontend/OpenMP/OMP.cpp b/llvm/lib/Frontend/OpenMP/OMP.cpp
index c1556ff3c74d7..b54cc90a14d83 100644
--- a/llvm/lib/Frontend/OpenMP/OMP.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMP.cpp
@@ -10,13 +10,19 @@
#include "llvm/ADT/ArrayRef.h"
#include "llvm/ADT/STLExtras.h"
+#include "llvm/ADT/SmallString.h"
#include "llvm/ADT/SmallVector.h"
#include "llvm/ADT/StringRef.h"
#include "llvm/ADT/StringSwitch.h"
+#include "llvm/Demangle/Demangle.h"
+#include "llvm/Frontend/OpenMP/OMPIRBuilder.h"
#include "llvm/Support/ErrorHandling.h"
+#include "llvm/Support/StringSaver.h"
#include <algorithm>
+#include <cstdio>
#include <iterator>
+#include <string>
#include <type_traits>
using namespace llvm;
@@ -186,4 +192,55 @@ bool isCombinedConstruct(Directive D) {
// Otherwise directive-name is a combined construct.
return !getLeafConstructs(D).empty() && !isCompositeConstruct(D);
}
+
+std::string prettityFunctionName(StringRef FunctionName) {
+ // Internalized functions have the right name, but simply a suffix.
+ if (FunctionName.ends_with(".internalized"))
+ return FunctionName.drop_back(sizeof("internalized")).str() +
+ " (internalized)";
+ unsigned LineNo = 0;
+ auto ParentName = deconstructOpenMPKernelName(FunctionName, LineNo);
+ if (LineNo == 0)
+ return FunctionName.str();
+ return ("omp target in " + ParentName + " @ " + std::to_string(LineNo) +
+ " (" + FunctionName + ")")
+ .str();
+}
+
+std::string deconstructOpenMPKernelName(StringRef KernelName,
+ unsigned &LineNo) {
+
+ // Only handle functions with an OpenMP kernel prefix for now. Naming scheme:
+ // __omp_offloading_<hex_hash1>_<hex_hash2>_<name>_l<line>_[<count>_]<suffix>
+ if (!KernelName.starts_with(TargetRegionEntryInfo::KernelNamePrefix))
+ return "";
+ auto SkipAfterNext = [](StringRef S, char Tgt, int &Remaining) {
+ return S.drop_while([&](char C) {
+ if (!Remaining)
+ return false;
+ Remaining -= (C == Tgt);
+ return true;
+ });
+ };
+ auto PrettyName = KernelName.drop_front(
+ sizeof(TargetRegionEntryInfo::KernelNamePrefix) - /*'\0'*/ 1);
+ int Remaining = 3;
+ PrettyName = SkipAfterNext(PrettyName, '_', Remaining);
+ if (Remaining)
+ return "";
+
+ // Look for the last '_l<line>'.
+ size_t LineIdx = PrettyName.find("_l");
+ if (LineIdx == StringRef::npos)
+ return "";
+ while (true) {
+ size_t NewLineIdx = PrettyName.find("_l", LineIdx + 2);
+ if (NewLineIdx == StringRef::npos)
+ break;
+ LineIdx = NewLineIdx;
+ }
+ if (PrettyName.drop_front(LineIdx + 2).consumeInteger(10, LineNo))
+ return "";
+ return demangle(PrettyName.take_front(LineIdx));
+}
} // namespace llvm::omp
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 77e350e7276ab..3f8e64315849e 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -8581,7 +8581,7 @@ void TargetRegionEntryInfo::getTargetRegionEntryFnName(
SmallVectorImpl<char> &Name, StringRef ParentName, unsigned DeviceID,
unsigned FileID, unsigned Line, unsigned Count) {
raw_svector_ostream OS(Name);
- OS << "__omp_offloading" << llvm::format("_%x", DeviceID)
+ OS << KernelNamePrefix << llvm::format("%x", DeviceID)
<< llvm::format("_%x_", FileID) << ParentName << "_l" << Line;
if (Count)
OS << "_" << Count;
diff --git a/offload/plugins-nextgen/common/include/ErrorReporting.h b/offload/plugins-nextgen/common/include/ErrorReporting.h
index 72cfb5273ae3c..bca7b27304a0b 100644
--- a/offload/plugins-nextgen/common/include/ErrorReporting.h
+++ b/offload/plugins-nextgen/common/include/ErrorReporting.h
@@ -17,6 +17,7 @@
#include "llvm/ADT/STLExtras.h"
#include "llvm/ADT/SmallString.h"
#include "llvm/ADT/StringRef.h"
+#include "llvm/Frontend/OpenMP/OMP.h"
#include "llvm/Support/ErrorHandling.h"
#include "llvm/Support/WithColor.h"
#include "llvm/Support/raw_ostream.h"
@@ -237,8 +238,11 @@ class ErrorReporter {
}
auto KTI = KTIR.getKernelTraceInfo(Idx);
- if (KTI.AsyncInfo && (AsyncInfoWrapperMatcher(*KTI.AsyncInfo)))
- reportError("Kernel '%s'", KTI.Kernel->getName());
+ if (KTI.AsyncInfo && (AsyncInfoWrapperMatcher(*KTI.AsyncInfo))) {
+ auto PrettyKernelName =
+ llvm::omp::prettityFunctionName(KTI.Kernel->getName());
+ reportError("Kernel '%s'", PrettyKernelName.c_str());
+ }
reportError("execution interrupted by hardware trap instruction");
if (KTI.AsyncInfo && (AsyncInfoWrapperMatcher(*KTI.AsyncInfo))) {
if (!KTI.LaunchTrace.empty())
@@ -284,10 +288,13 @@ class ErrorReporter {
for (uint32_t Idx = 0, I = 0; I < NumKTIs; ++Idx) {
auto KTI = KTIR.getKernelTraceInfo(Idx);
+ auto PrettyKernelName =
+ llvm::omp::prettityFunctionName(KTI.Kernel->getName());
if (NumKTIs == 1)
- print(BoldLightPurple, "Kernel '%s'\n", KTI.Kernel->getName());
+ print(BoldLightPurple, "Kernel '%s'\n", PrettyKernelName.c_str());
else
- print(BoldLightPurple, "Kernel %d: '%s'\n", I, KTI.Kernel->getName());
+ print(BoldLightPurple, "Kernel %d: '%s'\n", I,
+ PrettyKernelName.c_str());
reportStackTrace(KTI.LaunchTrace);
++I;
}
diff --git a/offload/src/CMakeLists.txt b/offload/src/CMakeLists.txt
index efa5cdab33ec9..344069b6fcdcf 100644
--- a/offload/src/CMakeLists.txt
+++ b/offload/src/CMakeLists.txt
@@ -28,6 +28,7 @@ add_llvm_library(omptarget
${LIBOMPTARGET_BINARY_INCLUDE_DIR}
LINK_COMPONENTS
+ FrontendOpenMP
Support
Object
diff --git a/offload/test/sanitizer/kernel_crash.c b/offload/test/sanitizer/kernel_crash.c
index 457d953a33a05..c69219d97d3d0 100644
--- a/offload/test/sanitizer/kernel_crash.c
+++ b/offload/test/sanitizer/kernel_crash.c
@@ -36,12 +36,12 @@ int main(void) {
}
}
// TRACE: Display 1 of the 3 last kernel launch traces
-// TRACE: Kernel 0: '__omp_offloading_{{.*}}_main_l30'
+// TRACE: Kernel 0: {{.*}} (__omp_offloading_{{.*}}_main_l30)
// TRACE: launchKernel
// NDEBG: main
// DEBUG: main {{.*}}kernel_crash.c:30
//
// CHECK: Display last 3 kernels launched:
-// CHECK: Kernel 0: '__omp_offloading_{{.*}}_main_l30'
-// CHECK: Kernel 1: '__omp_offloading_{{.*}}_main_l27'
-// CHECK: Kernel 2: '__omp_offloading_{{.*}}_main_l24'
+// CHECK: Kernel 0: {{.*}} (__omp_offloading_{{.*}}_main_l30)
+// CHECK: Kernel 1: {{.*}} (__omp_offloading_{{.*}}_main_l27)
+// CHECK: Kernel 2: {{.*}} (__omp_offloading_{{.*}}_main_l24)
diff --git a/offload/test/sanitizer/kernel_crash_async.c b/offload/test/sanitizer/kernel_crash_async.c
index 6aebf1b42a535..6a0461b0045b2 100644
--- a/offload/test/sanitizer/kernel_crash_async.c
+++ b/offload/test/sanitizer/kernel_crash_async.c
@@ -34,7 +34,7 @@ int main(void) {
#pragma omp taskwait
}
-// TRACE: Kernel {{.*}}'__omp_offloading_{{.*}}_main_
+// TRACE: Kernel {{.*}} (__omp_offloading_{{.*}}_main_l30)
// TRACE: launchKernel
//
-// CHECK-DAG: Kernel {{[0-9]}}: '__omp_offloading_{{.*}}_main_l30'
+// CHECK: Kernel {{[0-9]}}: {{.*}} (__omp_offloading_{{.*}}_main_l30)
diff --git a/offload/test/sanitizer/kernel_crash_many.c b/offload/test/sanitizer/kernel_crash_many.c
index 9e3f4f1630acd..25986e0a459c1 100644
--- a/offload/test/sanitizer/kernel_crash_many.c
+++ b/offload/test/sanitizer/kernel_crash_many.c
@@ -30,42 +30,42 @@ int main(void) {
}
}
// CHECK: Display 8 of the 8 last kernel launch traces
-// CHECK: Kernel 0: '__omp_offloading_{{.*}}_main_l27'
+// CHECK: Kernel 0: {{.*}} (__omp_offloading_{{.*}}_main_l27)
// CHECK: launchKernel
// NDEBG: main
// DEBUG: main {{.*}}kernel_crash_many.c:27
//
-// CHECK: Kernel 1: '__omp_offloading_{{.*}}_main_l23'
+// CHECK: Kernel 1: {{.*}} (__omp_offloading_{{.*}}_main_l23)
// CHECK: launchKernel
// NDEBG: main
// DEBUG: main {{.*}}kernel_crash_many.c:
//
-// CHECK: Kernel 2: '__omp_offloading_{{.*}}_main_l23'
+// CHECK: Kernel 2: {{.*}} (__omp_offloading_{{.*}}_main_l23)
// CHECK: launchKernel
// NDEBG: main
// DEBUG: main {{.*}}kernel_crash_many.c:
//
-// CHECK: Kernel 3: '__omp_offloading_{{.*}}_main_l23'
+// CHECK: Kernel 3: {{.*}} (__omp_offloading_{{.*}}_main_l23)
// CHECK: launchKernel
// NDEBG: main
// DEBUG: main {{.*}}kernel_crash_many.c:
//
-// CHECK: Kernel 4: '__omp_offloading_{{.*}}_main_l23'
+// CHECK: Kernel 4: {{.*}} (__omp_offloading_{{.*}}_main_l23)
// CHECK: launchKernel
// NDEBG: main
// DEBUG: main {{.*}}kernel_crash_many.c:
//
-// CHECK: Kernel 5: '__omp_offloading_{{.*}}_main_l23'
+// CHECK: Kernel 5: {{.*}} (__omp_offloading_{{.*}}_main_l23)
// CHECK: launchKernel
// NDEBG: main
// DEBUG: main {{.*}}kernel_crash_many.c:
//
-// CHECK: Kernel 6: '__omp_offloading_{{.*}}_main_l23'
+// CHECK: Kernel 6: {{.*}} (__omp_offloading_{{.*}}_main_l23)
// CHECK: launchKernel
// NDEBG: main
// DEBUG: main {{.*}}kernel_crash_many.c:
//
-// CHECK: Kernel 7: '__omp_offloading_{{.*}}_main_l23'
+// CHECK: Kernel 7: {{.*}} (__omp_offloading_{{.*}}_main_l23)
// CHECK: launchKernel
// NDEBG: main
// DEBUG: main {{.*}}kernel_crash_many.c:
diff --git a/offload/test/sanitizer/kernel_crash_single.c b/offload/test/sanitizer/kernel_crash_single.c
index 16a8159f074e5..075c3de7ffabb 100644
--- a/offload/test/sanitizer/kernel_crash_single.c
+++ b/offload/test/sanitizer/kernel_crash_single.c
@@ -27,10 +27,10 @@ int main(void) {
}
}
// TRACE: Display kernel launch trace
-// TRACE: Kernel '__omp_offloading_{{.*}}_main_l24'
+// TRACE: Kernel {{.*}} (__omp_offloading_{{.*}}_main_l24)
// TRACE: launchKernel
// NDEBG: main
// DEBUG: main {{.*}}kernel_crash_single.c:24
//
// CHECK: Display only launched kernel:
-// CHECK: Kernel '__omp_offloading_{{.*}}_main_l24'
+// CHECK: Kernel {{.*}} (__omp_offloading_{{.*}}_main_l24)
diff --git a/offload/test/sanitizer/kernel_trap.c b/offload/test/sanitizer/kernel_trap.c
index 13fe6f2fb71e8..db243001c9056 100644
--- a/offload/test/sanitizer/kernel_trap.c
+++ b/offload/test/sanitizer/kernel_trap.c
@@ -35,8 +35,10 @@ int main(void) {
{
}
}
-// CHECK: OFFLOAD ERROR: Kernel '__omp_offloading_{{.*}}_main_l30'
+// clang-format off
+// CHECK: OFFLOAD ERROR: Kernel 'omp target in main @ 30 (__omp_offloading_{{.*}}_main_l30)'
// CHECK: OFFLOAD ERROR: execution interrupted by hardware trap instruction
// TRACE: launchKernel
// NDEBG: main
// DEBUG: main {{.*}}kernel_trap.c:
+// clang-format on
diff --git a/offload/test/sanitizer/kernel_trap_async.c b/offload/test/sanitizer/kernel_trap_async.c
index 65e8880798343..ee0d772fef9b8 100644
--- a/offload/test/sanitizer/kernel_trap_async.c
+++ b/offload/test/sanitizer/kernel_trap_async.c
@@ -34,7 +34,9 @@ int main(void) {
#pragma omp taskwait
}
-// CHECK: OFFLOAD ERROR: Kernel '__omp_offloading_{{.*}}_main_l30'
+// clang-format off
+// CHECK: OFFLOAD ERROR: Kernel {{.*}} (__omp_offloading_{{.*}}_main_l30)
// CHECK: OFFLOAD ERROR: execution interrupted by hardware trap instruction
// TRACE: launchKernel
// DEBUG: kernel_trap_async.c:
+// clang-format on
diff --git a/offload/test/sanitizer/kernel_trap_many.c b/offload/test/sanitizer/kernel_trap_many.c
index 3f1796e8913ea..b3bdad9f07b4a 100644
--- a/offload/test/sanitizer/kernel_trap_many.c
+++ b/offload/test/sanitizer/kernel_trap_many.c
@@ -29,7 +29,7 @@ int main(void) {
__builtin_trap();
}
}
-// TRACE: OFFLOAD ERROR: Kernel '__omp_offloading_{{.*}}_main_l27'
+// TRACE: OFFLOAD ERROR: Kernel {{.*}} (__omp_offloading_{{.*}}_main_l27)
// TRACE: OFFLOAD ERROR: execution interrupted by hardware trap instruction
// TRACE: launchKernel
// NDEBG: main
|
jdoerfert
force-pushed
the
pr/kernel_demangling
branch
from
0a2552c to
4a0ba06
Compare
jhuber6
reviewed
Jul 31, 2024
You can test this locally with the following command:git-clang-format --diff 9a1013220b668d846e63f241203b80515dee0a03 fb4214bef997b8665b61e275c6e3e6851237f912 --extensions cpp,c,h -- offload/test/sanitizer/kernel_trap.cpp llvm/include/llvm/Frontend/OpenMP/OMP.h llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h llvm/lib/Frontend/OpenMP/OMP.cpp llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp offload/plugins-nextgen/common/include/ErrorReporting.h offload/test/sanitizer/kernel_crash.c offload/test/sanitizer/kernel_crash_async.c offload/test/sanitizer/kernel_crash_many.c offload/test/sanitizer/kernel_crash_single.c offload/test/sanitizer/kernel_trap.c offload/test/sanitizer/kernel_trap_async.c offload/test/sanitizer/kernel_trap_many.cView the diff from clang-format here.diff --git a/offload/test/sanitizer/kernel_trap.cpp b/offload/test/sanitizer/kernel_trap.cpp
index 899b608d57..4b454ce782 100644
--- a/offload/test/sanitizer/kernel_trap.cpp
+++ b/offload/test/sanitizer/kernel_trap.cpp
@@ -22,18 +22,13 @@ struct S {};
template <typename T> void cxx_function_name(int I, T *) {
#pragma omp target
- {
- }
+ {}
#pragma omp target
- {
- }
+ {}
#pragma omp target
- {
- __builtin_trap();
- }
+ { __builtin_trap(); }
#pragma omp target
- {
- }
+ {}
}
int main(void) {
|
jdoerfert
force-pushed
the
pr/kernel_demangling
branch
2 times, most recently
from
c2269ac to
b613d12
Compare
jhuber6
reviewed
Jul 31, 2024
jdoerfert
force-pushed
the
pr/kernel_demangling
branch
from
b613d12 to
67bcaba
Compare
…name The kernel names for OpenMP are manually mangled and not ideal when we report something to the user. We demangle them now, providing the function and line number of the target region, together with the actual kernel name.
jdoerfert
force-pushed
the
pr/kernel_demangling
branch
from
67bcaba to
fb4214b
Compare
jhuber6
approved these changes
Aug 1, 2024
|
I ran into a similar issue downstream and extended |
We could possibly make the OpenMP offloading kernel mangling compatible with the mangling format, then we'd get it for free. The issue is that the manged kernel name needs to contain something pseudo-unique (We use the file's inode + the line number currently). Everything else like |
We could make it regular mangling, but I'm not sold it is worth it. @dwoodwor-intel We should decide if we want to use Itanium mangling or not, if not, I'd be happy to review your changes and you can remove some of the stuff in this PR. We can always add the original kernel name, which is useful since other tools emit it, e.g., rocprof. |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
The kernel names for OpenMP are manually mangled and not ideal when we report something to the user. We demangle them now, providing the function and line number of the target region, together with the actual kernel name.