[Offload][OpenMP] Prettify error messages by "demangling" the kernel name (#101400)

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.

Author: jdoerfert

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