[Offload][OpenMP] Prettify error messages by "demangling" the kernel name

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 prettifyFunctionName(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

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;

llvm/lib/Frontend/OpenMP/CMakeLists.txt

@@ -17,6 +17,7 @@ add_llvm_component_library(LLVMFrontendOpenMP
   TargetParser
   TransformUtils
   Analysis
+  Demangle
   MC
   Scalar
   BitReader

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,42 @@ bool isCombinedConstruct(Directive D) {
   // Otherwise directive-name is a combined construct.
   return !getLeafConstructs(D).empty() && !isCompositeConstruct(D);
 }
+
+std::string prettifyFunctionName(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 PrettyName = KernelName.drop_front(
+      sizeof(TargetRegionEntryInfo::KernelNamePrefix) - /*'\0'*/ 1);
+  for (int I = 0; I < 3; ++I) {
+    PrettyName = PrettyName.drop_while([](char c) { return c != '_'; });
+    PrettyName = PrettyName.drop_front();
+  }
+
+  // Look for the last '_l<line>'.
+  size_t LineIdx = PrettyName.rfind("_l");
+  if (LineIdx == StringRef::npos)
+    return "";
+  if (PrettyName.drop_front(LineIdx + 2).consumeInteger(10, LineNo))
+    return "";
+  return demangle(PrettyName.take_front(LineIdx));
+}
 } // namespace llvm::omp

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;

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::prettifyFunctionName(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::prettifyFunctionName(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;
     }

offload/src/CMakeLists.txt

@@ -28,6 +28,7 @@ add_llvm_library(omptarget
   ${LIBOMPTARGET_BINARY_INCLUDE_DIR}
 
   LINK_COMPONENTS
+  FrontendOpenMP
   Support
   Object
 

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)

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)

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:

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)

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

offload/test/sanitizer/kernel_trap.cpp

@@ -0,0 +1,52 @@
+
+// clang-format off
+// RUN: %libomptarget-compilexx-generic
+// RUN: %not --crash env -u LLVM_DISABLE_SYMBOLIZATION OFFLOAD_TRACK_NUM_KERNEL_LAUNCH_TRACES=1 %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK,TRACE,NDEBG 
+// RUN: %not --crash %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK
+// RUN: %libomptarget-compilexx-generic -g
+// RUN: %not --crash env -u LLVM_DISABLE_SYMBOLIZATION OFFLOAD_TRACK_NUM_KERNEL_LAUNCH_TRACES=1 %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK,TRACE,DEBUG
+// RUN: %not --crash %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK
+// clang-format on
+
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: aarch64-unknown-linux-gnu-LTO
+// UNSUPPORTED: x86_64-pc-linux-gnu
+// UNSUPPORTED: x86_64-pc-linux-gnu-LTO
+// UNSUPPORTED: s390x-ibm-linux-gnu
+// UNSUPPORTED: s390x-ibm-linux-gnu-LTO
+
+struct S {};
+
+template <typename T> void cxx_function_name(int I, T *) {
+
+#pragma omp target
+  {
+  }
+#pragma omp target
+  {
+  }
+#pragma omp target
+  {
+    __builtin_trap();
+  }
+#pragma omp target
+  {
+  }
+}
+
+int main(void) {
+  struct S s;
+  cxx_function_name(1, &s);
+}
+
+// clang-format off
+// CHECK: OFFLOAD ERROR: Kernel 'omp target in void cxx_function_name<S>(int, S*) @ [[LINE:[0-9]+]] (__omp_offloading_{{.*}}__Z17cxx_function_nameI1SEviPT__l[[LINE]])'
+// CHECK: OFFLOAD ERROR: execution interrupted by hardware trap instruction
+// TRACE:     launchKernel
+// NDEBG:     cxx_function_name<S>(int, S*)
+// NDEBG:     main
+// DEBUG:     cxx_function_name<S>(int, S*) {{.*}}kernel_trap.cpp:
+// DEBUG:     main {{.*}}kernel_trap.cpp:
+// clang-format on

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

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