[Openmp-commits] [llvm] [openmp] [Offload] Allow to record kernel launch stack traces (PR #100472)
Johannes Doerfert via Openmp-commits
openmp-commits at lists.llvm.org
Wed Jul 31 11:49:17 PDT 2024
https://github.com/jdoerfert updated https://github.com/llvm/llvm-project/pull/100472
>From e0816f3d7482e38dce23e37ebc58c957733f8d38 Mon Sep 17 00:00:00 2001
From: Johannes Doerfert <johannes at jdoerfert.de>
Date: Tue, 23 Jul 2024 20:45:39 -0700
Subject: [PATCH] [Offload] Allow to record kernel launch stack traces
Similar to (de)allocation traces, we can record kernel launch stack
traces and display them in case of an error. However, the AMD GPU plugin
signal handler, which is invoked on memroy faults, cannot pinpoint the
offending kernel. Insteade print `<NUM>`, set via
`OFFLOAD_TRACK_NUM_KERNEL_LAUNCH_TRACES=<NUM>`, many traces. The
recoding/record uses a ring buffer of fixed size (for now 8).
For `trap` errors, we print the actual kernel name, and trace if
recorded.
---
offload/include/Shared/EnvironmentVar.h | 6 +-
.../plugins-nextgen/amdgpu/dynamic_hsa/hsa.h | 1 +
offload/plugins-nextgen/amdgpu/src/rtl.cpp | 72 +++++++++++++---
.../common/include/ErrorReporting.h | 85 +++++++++++++++++++
.../common/include/PluginInterface.h | 46 ++++++++++
.../common/src/PluginInterface.cpp | 12 +++
offload/test/sanitizer/kernel_crash.c | 47 ++++++++++
offload/test/sanitizer/kernel_crash_async.c | 40 +++++++++
offload/test/sanitizer/kernel_crash_many.c | 73 ++++++++++++++++
offload/test/sanitizer/kernel_crash_single.c | 36 ++++++++
offload/test/sanitizer/kernel_trap.c | 42 +++++++++
offload/test/sanitizer/kernel_trap_async.c | 40 +++++++++
offload/test/sanitizer/kernel_trap_many.c | 36 ++++++++
openmp/docs/design/Runtimes.rst | 7 ++
14 files changed, 531 insertions(+), 12 deletions(-)
create mode 100644 offload/test/sanitizer/kernel_crash.c
create mode 100644 offload/test/sanitizer/kernel_crash_async.c
create mode 100644 offload/test/sanitizer/kernel_crash_many.c
create mode 100644 offload/test/sanitizer/kernel_crash_single.c
create mode 100644 offload/test/sanitizer/kernel_trap.c
create mode 100644 offload/test/sanitizer/kernel_trap_async.c
create mode 100644 offload/test/sanitizer/kernel_trap_many.c
diff --git a/offload/include/Shared/EnvironmentVar.h b/offload/include/Shared/EnvironmentVar.h
index 4cbdad695a0ee..82f434e91a85b 100644
--- a/offload/include/Shared/EnvironmentVar.h
+++ b/offload/include/Shared/EnvironmentVar.h
@@ -28,6 +28,7 @@ struct StringParser {
/// Class for reading and checking environment variables. Currently working with
/// integer, floats, std::string and bool types.
template <typename Ty> class Envar {
+ llvm::StringRef Name;
Ty Data;
bool IsPresent;
bool Initialized;
@@ -53,7 +54,7 @@ template <typename Ty> class Envar {
/// take the value read from the environment variable, or the default if it
/// was not set or not correct. This constructor is not fallible.
Envar(llvm::StringRef Name, Ty Default = Ty())
- : Data(Default), IsPresent(false), Initialized(true) {
+ : Name(Name), Data(Default), IsPresent(false), Initialized(true) {
if (const char *EnvStr = getenv(Name.data())) {
// Check whether the envar is defined and valid.
@@ -84,6 +85,9 @@ template <typename Ty> class Envar {
/// Get the definitive value.
operator Ty() const { return get(); }
+ /// Return the environment variable name.
+ llvm::StringRef getName() const { return Name; }
+
/// Indicate whether the environment variable was defined and valid.
bool isPresent() const { return IsPresent; }
diff --git a/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.h b/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.h
index 64a1d3308aed0..5d9fb5d7dc7cd 100644
--- a/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.h
+++ b/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.h
@@ -31,6 +31,7 @@ typedef enum {
HSA_STATUS_ERROR = 0x1000,
HSA_STATUS_ERROR_INVALID_CODE_OBJECT = 0x1010,
HSA_STATUS_ERROR_NOT_INITIALIZED = 0x100B,
+ HSA_STATUS_ERROR_EXCEPTION = 0x1016,
} hsa_status_t;
hsa_status_t hsa_status_string(hsa_status_t status, const char **status_string);
diff --git a/offload/plugins-nextgen/amdgpu/src/rtl.cpp b/offload/plugins-nextgen/amdgpu/src/rtl.cpp
index e6643d3260eb4..604683370cd27 100644
--- a/offload/plugins-nextgen/amdgpu/src/rtl.cpp
+++ b/offload/plugins-nextgen/amdgpu/src/rtl.cpp
@@ -13,13 +13,16 @@
#include <atomic>
#include <cassert>
#include <cstddef>
+#include <cstdint>
#include <deque>
+#include <functional>
#include <mutex>
#include <string>
#include <system_error>
#include <unistd.h>
#include <unordered_map>
+#include "ErrorReporting.h"
#include "Shared/APITypes.h"
#include "Shared/Debug.h"
#include "Shared/Environment.h"
@@ -43,6 +46,7 @@
#include "llvm/Support/FileSystem.h"
#include "llvm/Support/MemoryBuffer.h"
#include "llvm/Support/Program.h"
+#include "llvm/Support/Signals.h"
#include "llvm/Support/raw_ostream.h"
#if !defined(__BYTE_ORDER__) || !defined(__ORDER_LITTLE_ENDIAN__) || \
@@ -685,12 +689,12 @@ struct AMDGPUQueueTy {
AMDGPUQueueTy() : Queue(nullptr), Mutex(), NumUsers(0) {}
/// Lazily initialize a new queue belonging to a specific agent.
- Error init(hsa_agent_t Agent, int32_t QueueSize) {
+ Error init(GenericDeviceTy &Device, hsa_agent_t Agent, int32_t QueueSize) {
if (Queue)
return Plugin::success();
hsa_status_t Status =
hsa_queue_create(Agent, QueueSize, HSA_QUEUE_TYPE_MULTI, callbackError,
- nullptr, UINT32_MAX, UINT32_MAX, &Queue);
+ &Device, UINT32_MAX, UINT32_MAX, &Queue);
return Plugin::check(Status, "Error in hsa_queue_create: %s");
}
@@ -875,10 +879,8 @@ struct AMDGPUQueueTy {
}
/// Callack that will be called when an error is detected on the HSA queue.
- static void callbackError(hsa_status_t Status, hsa_queue_t *Source, void *) {
- auto Err = Plugin::check(Status, "Received error in queue %p: %s", Source);
- FATAL_MESSAGE(1, "%s", toString(std::move(Err)).data());
- }
+ static void callbackError(hsa_status_t Status, hsa_queue_t *Source,
+ void *Data);
/// The HSA queue.
hsa_queue_t *Queue;
@@ -1484,6 +1486,8 @@ struct AMDGPUStreamTy {
return true;
}
+ const AMDGPUQueueTy *getQueue() const { return Queue; }
+
/// Record the state of the stream on an event.
Error recordEvent(AMDGPUEventTy &Event) const;
@@ -1594,7 +1598,7 @@ struct AMDGPUStreamManagerTy final
using ResourcePoolTy = GenericDeviceResourceManagerTy<ResourceRef>;
AMDGPUStreamManagerTy(GenericDeviceTy &Device, hsa_agent_t HSAAgent)
- : GenericDeviceResourceManagerTy(Device),
+ : GenericDeviceResourceManagerTy(Device), Device(Device),
OMPX_QueueTracking("LIBOMPTARGET_AMDGPU_HSA_QUEUE_BUSY_TRACKING", true),
NextQueue(0), Agent(HSAAgent) {}
@@ -1603,7 +1607,7 @@ struct AMDGPUStreamManagerTy final
QueueSize = HSAQueueSize;
MaxNumQueues = NumHSAQueues;
// Initialize one queue eagerly
- if (auto Err = Queues.front().init(Agent, QueueSize))
+ if (auto Err = Queues.front().init(Device, Agent, QueueSize))
return Err;
return GenericDeviceResourceManagerTy::init(InitialSize);
@@ -1660,7 +1664,7 @@ struct AMDGPUStreamManagerTy final
}
// Make sure the queue is initialized, then add user & assign.
- if (auto Err = Queues[Index].init(Agent, QueueSize))
+ if (auto Err = Queues[Index].init(Device, Agent, QueueSize))
return Err;
Queues[Index].addUser();
Stream->Queue = &Queues[Index];
@@ -1668,6 +1672,9 @@ struct AMDGPUStreamManagerTy final
return Plugin::success();
}
+ /// The device associated with this stream.
+ GenericDeviceTy &Device;
+
/// Envar for controlling the tracking of busy HSA queues.
BoolEnvar OMPX_QueueTracking;
@@ -3074,7 +3081,7 @@ struct AMDGPUPluginTy final : public GenericPluginTy {
Initialized = true;
// Register event handler to detect memory errors on the devices.
- Status = hsa_amd_register_system_event_handler(eventHandler, nullptr);
+ Status = hsa_amd_register_system_event_handler(eventHandler, this);
if (auto Err = Plugin::check(
Status, "Error in hsa_amd_register_system_event_handler: %s"))
return std::move(Err);
@@ -3209,7 +3216,8 @@ struct AMDGPUPluginTy final : public GenericPluginTy {
private:
/// Event handler that will be called by ROCr if an event is detected.
- static hsa_status_t eventHandler(const hsa_amd_event_t *Event, void *) {
+ static hsa_status_t eventHandler(const hsa_amd_event_t *Event,
+ void *PluginPtr) {
if (Event->event_type != HSA_AMD_GPU_MEMORY_FAULT_EVENT)
return HSA_STATUS_SUCCESS;
@@ -3240,6 +3248,26 @@ struct AMDGPUPluginTy final : public GenericPluginTy {
uint32_t Node = -1;
hsa_agent_get_info(Event->memory_fault.agent, HSA_AGENT_INFO_NODE, &Node);
+ AMDGPUPluginTy &Plugin = *reinterpret_cast<AMDGPUPluginTy *>(PluginPtr);
+ for (uint32_t I = 0, E = Plugin.getNumDevices();
+ Node != uint32_t(-1) && I < E; ++I) {
+ AMDGPUDeviceTy &AMDGPUDevice =
+ reinterpret_cast<AMDGPUDeviceTy &>(Plugin.getDevice(I));
+ auto KernelTraceInfoRecord =
+ AMDGPUDevice.KernelLaunchTraces.getExclusiveAccessor();
+
+ uint32_t DeviceNode = -1;
+ if (auto Err =
+ AMDGPUDevice.getDeviceAttr(HSA_AGENT_INFO_NODE, DeviceNode)) {
+ consumeError(std::move(Err));
+ continue;
+ }
+ if (DeviceNode != Node)
+ continue;
+
+ ErrorReporter::reportKernelTraces(AMDGPUDevice, *KernelTraceInfoRecord);
+ }
+
// Abort the execution since we do not recover from this error.
FATAL_MESSAGE(1,
"Memory access fault by GPU %" PRIu32 " (agent 0x%" PRIx64
@@ -3480,6 +3508,28 @@ void *AMDGPUDeviceTy::allocate(size_t Size, void *, TargetAllocTy Kind) {
return Alloc;
}
+void AMDGPUQueueTy::callbackError(hsa_status_t Status, hsa_queue_t *Source,
+ void *Data) {
+ auto &AMDGPUDevice = *reinterpret_cast<AMDGPUDeviceTy *>(Data);
+
+ if (Status == HSA_STATUS_ERROR_EXCEPTION) {
+ auto KernelTraceInfoRecord =
+ AMDGPUDevice.KernelLaunchTraces.getExclusiveAccessor();
+ std::function<bool(__tgt_async_info &)> AsyncInfoWrapperMatcher =
+ [=](__tgt_async_info &AsyncInfo) {
+ auto *Stream = reinterpret_cast<AMDGPUStreamTy *>(AsyncInfo.Queue);
+ if (!Stream || !Stream->getQueue())
+ return false;
+ return Stream->getQueue()->Queue == Source;
+ };
+ ErrorReporter::reportTrapInKernel(AMDGPUDevice, *KernelTraceInfoRecord,
+ AsyncInfoWrapperMatcher);
+ }
+
+ auto Err = Plugin::check(Status, "Received error in queue %p: %s", Source);
+ FATAL_MESSAGE(1, "%s", toString(std::move(Err)).data());
+}
+
} // namespace plugin
} // namespace target
} // namespace omp
diff --git a/offload/plugins-nextgen/common/include/ErrorReporting.h b/offload/plugins-nextgen/common/include/ErrorReporting.h
index a8efaad35bb10..72cfb5273ae3c 100644
--- a/offload/plugins-nextgen/common/include/ErrorReporting.h
+++ b/offload/plugins-nextgen/common/include/ErrorReporting.h
@@ -14,6 +14,7 @@
#include "PluginInterface.h"
#include "Shared/EnvironmentVar.h"
+#include "llvm/ADT/STLExtras.h"
#include "llvm/ADT/SmallString.h"
#include "llvm/ADT/StringRef.h"
#include "llvm/Support/ErrorHandling.h"
@@ -216,6 +217,90 @@ class ErrorReporter {
getAllocTyName(ATI->Kind).data(), DevicePtr);
#undef DEALLOCATION_ERROR
}
+
+ /// Report that a kernel encountered a trap instruction.
+ static void reportTrapInKernel(
+ GenericDeviceTy &Device, KernelTraceInfoRecordTy &KTIR,
+ std::function<bool(__tgt_async_info &)> AsyncInfoWrapperMatcher) {
+ assert(AsyncInfoWrapperMatcher && "A matcher is required");
+
+ uint32_t Idx = 0;
+ for (uint32_t I = 0, E = KTIR.size(); I < E; ++I) {
+ auto KTI = KTIR.getKernelTraceInfo(I);
+ if (KTI.Kernel == nullptr)
+ break;
+ // Skip kernels issued in other queues.
+ if (KTI.AsyncInfo && !(AsyncInfoWrapperMatcher(*KTI.AsyncInfo)))
+ continue;
+ Idx = I;
+ break;
+ }
+
+ auto KTI = KTIR.getKernelTraceInfo(Idx);
+ if (KTI.AsyncInfo && (AsyncInfoWrapperMatcher(*KTI.AsyncInfo)))
+ reportError("Kernel '%s'", KTI.Kernel->getName());
+ reportError("execution interrupted by hardware trap instruction");
+ if (KTI.AsyncInfo && (AsyncInfoWrapperMatcher(*KTI.AsyncInfo))) {
+ if (!KTI.LaunchTrace.empty())
+ reportStackTrace(KTI.LaunchTrace);
+ else
+ print(Yellow, "Use '%s=1' to show the stack trace of the kernel\n",
+ Device.OMPX_TrackNumKernelLaunches.getName().data());
+ }
+ abort();
+ }
+
+ /// Report the kernel traces taken from \p KTIR, up to
+ /// OFFLOAD_TRACK_NUM_KERNEL_LAUNCH_TRACES many.
+ static void reportKernelTraces(GenericDeviceTy &Device,
+ KernelTraceInfoRecordTy &KTIR) {
+ uint32_t NumKTIs = 0;
+ for (uint32_t I = 0, E = KTIR.size(); I < E; ++I) {
+ auto KTI = KTIR.getKernelTraceInfo(I);
+ if (KTI.Kernel == nullptr)
+ break;
+ ++NumKTIs;
+ }
+ if (NumKTIs == 0) {
+ print(BoldRed, "No kernel launches known\n");
+ return;
+ }
+
+ uint32_t TracesToShow =
+ std::min(Device.OMPX_TrackNumKernelLaunches.get(), NumKTIs);
+ if (TracesToShow == 0) {
+ if (NumKTIs == 1)
+ print(BoldLightPurple, "Display only launched kernel:\n");
+ else
+ print(BoldLightPurple, "Display last %u kernels launched:\n", NumKTIs);
+ } else {
+ if (NumKTIs == 1)
+ print(BoldLightPurple, "Display kernel launch trace:\n");
+ else
+ print(BoldLightPurple,
+ "Display %u of the %u last kernel launch traces:\n", TracesToShow,
+ NumKTIs);
+ }
+
+ for (uint32_t Idx = 0, I = 0; I < NumKTIs; ++Idx) {
+ auto KTI = KTIR.getKernelTraceInfo(Idx);
+ if (NumKTIs == 1)
+ print(BoldLightPurple, "Kernel '%s'\n", KTI.Kernel->getName());
+ else
+ print(BoldLightPurple, "Kernel %d: '%s'\n", I, KTI.Kernel->getName());
+ reportStackTrace(KTI.LaunchTrace);
+ ++I;
+ }
+
+ if (NumKTIs != 1) {
+ print(Yellow,
+ "Use '%s=<num>' to adjust the number of shown stack traces (%u "
+ "now, up to %zu)\n",
+ Device.OMPX_TrackNumKernelLaunches.getName().data(),
+ Device.OMPX_TrackNumKernelLaunches.get(), KTIR.size());
+ }
+ // TODO: Let users know how to serialize kernels
+ }
};
} // namespace plugin
diff --git a/offload/plugins-nextgen/common/include/PluginInterface.h b/offload/plugins-nextgen/common/include/PluginInterface.h
index 49a3943a27429..81823338fe211 100644
--- a/offload/plugins-nextgen/common/include/PluginInterface.h
+++ b/offload/plugins-nextgen/common/include/PluginInterface.h
@@ -412,6 +412,44 @@ struct AllocationTraceInfoTy {
std::mutex Lock;
};
+/// Information about an allocation, when it has been allocated, and when/if it
+/// has been deallocated, for error reporting purposes.
+struct KernelTraceInfoTy {
+
+ /// The launched kernel.
+ GenericKernelTy *Kernel;
+
+ /// The stack trace of the launch itself.
+ std::string LaunchTrace;
+
+ /// The async info the kernel was launched in.
+ __tgt_async_info *AsyncInfo;
+};
+
+struct KernelTraceInfoRecordTy {
+ KernelTraceInfoRecordTy() { KTIs.fill({}); }
+
+ /// Return the (maximal) record size.
+ auto size() const { return KTIs.size(); }
+
+ /// Create a new kernel trace info and add it into the record.
+ void emplace(GenericKernelTy *Kernel, const std::string &&StackTrace,
+ __tgt_async_info *AsyncInfo) {
+ KTIs[Idx] = {Kernel, std::move(StackTrace), AsyncInfo};
+ Idx = (Idx + 1) % size();
+ }
+
+ /// Return the \p I'th last kernel trace info.
+ auto getKernelTraceInfo(int32_t I) const {
+ // Note that kernel trace infos "grow forward", so lookup is backwards.
+ return KTIs[(Idx - I - 1 + size()) % size()];
+ }
+
+private:
+ std::array<KernelTraceInfoTy, 8> KTIs;
+ unsigned Idx = 0;
+};
+
/// Class representing a map of host pinned allocations. We track these pinned
/// allocations, so memory tranfers invloving these buffers can be optimized.
class PinnedAllocationMapTy {
@@ -900,6 +938,14 @@ struct GenericDeviceTy : public DeviceAllocatorTy {
/// been deallocated, both for error reporting purposes.
ProtectedObj<DenseMap<void *, AllocationTraceInfoTy *>> AllocationTraces;
+ /// Map to record kernel have been launchedl, for error reporting purposes.
+ ProtectedObj<KernelTraceInfoRecordTy> KernelLaunchTraces;
+
+ /// Environment variable to determine if stack traces for kernel launches are
+ /// tracked.
+ UInt32Envar OMPX_TrackNumKernelLaunches =
+ UInt32Envar("OFFLOAD_TRACK_NUM_KERNEL_LAUNCH_TRACES", 0);
+
private:
/// Get and set the stack size and heap size for the device. If not used, the
/// plugin can implement the setters as no-op and setting the output
diff --git a/offload/plugins-nextgen/common/src/PluginInterface.cpp b/offload/plugins-nextgen/common/src/PluginInterface.cpp
index 3b6f545fe91fa..c3ecbcc62f71f 100644
--- a/offload/plugins-nextgen/common/src/PluginInterface.cpp
+++ b/offload/plugins-nextgen/common/src/PluginInterface.cpp
@@ -1468,6 +1468,18 @@ Error GenericDeviceTy::launchKernel(void *EntryPtr, void **ArgPtrs,
GenericKernelTy &GenericKernel =
*reinterpret_cast<GenericKernelTy *>(EntryPtr);
+ {
+ std::string StackTrace;
+ if (OMPX_TrackNumKernelLaunches) {
+ llvm::raw_string_ostream OS(StackTrace);
+ llvm::sys::PrintStackTrace(OS);
+ }
+
+ auto KernelTraceInfoRecord = KernelLaunchTraces.getExclusiveAccessor();
+ (*KernelTraceInfoRecord)
+ .emplace(&GenericKernel, std::move(StackTrace), AsyncInfo);
+ }
+
auto Err = GenericKernel.launch(*this, ArgPtrs, ArgOffsets, KernelArgs,
AsyncInfoWrapper);
diff --git a/offload/test/sanitizer/kernel_crash.c b/offload/test/sanitizer/kernel_crash.c
new file mode 100644
index 0000000000000..457d953a33a05
--- /dev/null
+++ b/offload/test/sanitizer/kernel_crash.c
@@ -0,0 +1,47 @@
+
+// clang-format off
+// RUN: %libomptarget-compile-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=TRACE,NDEBG
+// RUN: %not --crash %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK
+// RUN: %libomptarget-compile-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=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
+
+#include <omp.h>
+
+int main(void) {
+ int *A = 0;
+#pragma omp target
+ {
+ }
+#pragma omp target
+ {
+ }
+#pragma omp target
+ {
+ *A = 42;
+ }
+#pragma omp target
+ {
+ }
+}
+// TRACE: Display 1 of the 3 last kernel launch traces
+// 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'
diff --git a/offload/test/sanitizer/kernel_crash_async.c b/offload/test/sanitizer/kernel_crash_async.c
new file mode 100644
index 0000000000000..6aebf1b42a535
--- /dev/null
+++ b/offload/test/sanitizer/kernel_crash_async.c
@@ -0,0 +1,40 @@
+
+// clang-format off
+// RUN: %libomptarget-compileopt-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=TRACE
+// RUN: %not --crash %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK
+// RUN: %libomptarget-compileopt-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=TRACE
+// 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
+
+#include <omp.h>
+
+int main(void) {
+ int *A = 0;
+#pragma omp target nowait
+ {
+ }
+#pragma omp target nowait
+ {
+ }
+#pragma omp target nowait
+ {
+ *A = 42;
+ }
+#pragma omp taskwait
+}
+
+// TRACE: Kernel {{.*}}'__omp_offloading_{{.*}}_main_
+// TRACE: launchKernel
+//
+// CHECK-DAG: 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
new file mode 100644
index 0000000000000..9e3f4f1630acd
--- /dev/null
+++ b/offload/test/sanitizer/kernel_crash_many.c
@@ -0,0 +1,73 @@
+
+// clang-format off
+// RUN: %libomptarget-compile-generic
+// RUN: %not --crash env -u LLVM_DISABLE_SYMBOLIZATION OFFLOAD_TRACK_NUM_KERNEL_LAUNCH_TRACES=24 %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK,NDEBG
+// RUN: %libomptarget-compile-generic -g
+// RUN: %not --crash env -u LLVM_DISABLE_SYMBOLIZATION OFFLOAD_TRACK_NUM_KERNEL_LAUNCH_TRACES=16 %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK,DEBUG
+// 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
+
+#include <omp.h>
+
+int main(void) {
+ int *A = 0;
+ for (int i = 0; i < 10; ++i) {
+#pragma omp target
+ {
+ }
+ }
+#pragma omp target
+ {
+ *A = 42;
+ }
+}
+// CHECK: Display 8 of the 8 last kernel launch traces
+// 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: launchKernel
+// NDEBG: main
+// DEBUG: main {{.*}}kernel_crash_many.c:
+//
+// 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: launchKernel
+// NDEBG: main
+// DEBUG: main {{.*}}kernel_crash_many.c:
+//
+// 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: launchKernel
+// NDEBG: main
+// DEBUG: main {{.*}}kernel_crash_many.c:
+//
+// 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: launchKernel
+// NDEBG: main
+// DEBUG: main {{.*}}kernel_crash_many.c:
+//
+// CHECK-NOT: Kernel {{[[0-9]]+}}:
diff --git a/offload/test/sanitizer/kernel_crash_single.c b/offload/test/sanitizer/kernel_crash_single.c
new file mode 100644
index 0000000000000..16a8159f074e5
--- /dev/null
+++ b/offload/test/sanitizer/kernel_crash_single.c
@@ -0,0 +1,36 @@
+
+// clang-format off
+// RUN: %libomptarget-compile-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=TRACE,NDEBG
+// RUN: %not --crash %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK
+// RUN: %libomptarget-compile-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=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
+
+#include <omp.h>
+
+int main(void) {
+ int *A = 0;
+#pragma omp target
+ {
+ *A = 42;
+ }
+}
+// TRACE: Display kernel launch trace
+// 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'
diff --git a/offload/test/sanitizer/kernel_trap.c b/offload/test/sanitizer/kernel_trap.c
new file mode 100644
index 0000000000000..13fe6f2fb71e8
--- /dev/null
+++ b/offload/test/sanitizer/kernel_trap.c
@@ -0,0 +1,42 @@
+
+// clang-format off
+// RUN: %libomptarget-compile-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-compile-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
+
+#include <omp.h>
+
+int main(void) {
+
+#pragma omp target
+ {
+ }
+#pragma omp target
+ {
+ }
+#pragma omp target
+ {
+ __builtin_trap();
+ }
+#pragma omp target
+ {
+ }
+}
+// CHECK: OFFLOAD ERROR: Kernel '__omp_offloading_{{.*}}_main_l30'
+// CHECK: OFFLOAD ERROR: execution interrupted by hardware trap instruction
+// TRACE: launchKernel
+// NDEBG: main
+// DEBUG: main {{.*}}kernel_trap.c:
diff --git a/offload/test/sanitizer/kernel_trap_async.c b/offload/test/sanitizer/kernel_trap_async.c
new file mode 100644
index 0000000000000..65e8880798343
--- /dev/null
+++ b/offload/test/sanitizer/kernel_trap_async.c
@@ -0,0 +1,40 @@
+
+// clang-format off
+// RUN: %libomptarget-compileopt-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
+// RUN: %not --crash %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK
+// RUN: %libomptarget-compileopt-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
+
+#include <omp.h>
+
+int main(void) {
+
+#pragma omp target nowait
+ {
+ }
+#pragma omp target nowait
+ {
+ }
+#pragma omp target nowait
+ {
+ __builtin_trap();
+ }
+#pragma omp taskwait
+}
+
+// CHECK: OFFLOAD ERROR: Kernel '__omp_offloading_{{.*}}_main_l30'
+// CHECK: OFFLOAD ERROR: execution interrupted by hardware trap instruction
+// TRACE: launchKernel
+// DEBUG: kernel_trap_async.c:
diff --git a/offload/test/sanitizer/kernel_trap_many.c b/offload/test/sanitizer/kernel_trap_many.c
new file mode 100644
index 0000000000000..3f1796e8913ea
--- /dev/null
+++ b/offload/test/sanitizer/kernel_trap_many.c
@@ -0,0 +1,36 @@
+
+// clang-format off
+// RUN: %libomptarget-compile-generic
+// RUN: %not --crash env -u LLVM_DISABLE_SYMBOLIZATION OFFLOAD_TRACK_NUM_KERNEL_LAUNCH_TRACES=24 %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=TRACE,NDEBG
+// RUN: %libomptarget-compile-generic -g
+// RUN: %not --crash env -u LLVM_DISABLE_SYMBOLIZATION OFFLOAD_TRACK_NUM_KERNEL_LAUNCH_TRACES=16 %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=TRACE,DEBUG
+// 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
+
+#include <omp.h>
+
+int main(void) {
+
+ for (int i = 0; i < 10; ++i) {
+#pragma omp target
+ {
+ }
+ }
+#pragma omp target
+ {
+ __builtin_trap();
+ }
+}
+// TRACE: OFFLOAD ERROR: Kernel '__omp_offloading_{{.*}}_main_l27'
+// TRACE: OFFLOAD ERROR: execution interrupted by hardware trap instruction
+// TRACE: launchKernel
+// NDEBG: main
+// DEBUG: main {{.*}}kernel_trap_many.c:
diff --git a/openmp/docs/design/Runtimes.rst b/openmp/docs/design/Runtimes.rst
index 20e2fa715eaac..ed002c8cf0f80 100644
--- a/openmp/docs/design/Runtimes.rst
+++ b/openmp/docs/design/Runtimes.rst
@@ -744,6 +744,7 @@ variables is defined below.
* ``LIBOMPTARGET_MIN_THREADS_FOR_LOW_TRIP_COUNT=<Num> (default: 32)``
* ``LIBOMPTARGET_REUSE_BLOCKS_FOR_HIGH_TRIP_COUNT=[TRUE/FALSE] (default TRUE)``
* ``OFFLOAD_TRACK_ALLOCATION_TRACES=[TRUE/FALSE] (default FALSE)``
+ * ``OFFLOAD_TRACK_NUM_KERNEL_LAUNCH_TRACES=<Num> (default 0)``
LIBOMPTARGET_DEBUG
""""""""""""""""""
@@ -1178,6 +1179,12 @@ This environment variable determines if the stack traces of allocations and
deallocations are tracked to aid in error reporting, e.g., in case of
double-free.
+OFFLOAD_TRACK_KERNEL_LAUNCH_TRACES
+""""""""""""""""""""""""""""""""""
+
+This environment variable determines how manytstack traces of kernel launches
+are tracked to aid in error reporting, e.g., what asynchronous kernel failed.
+
.. _libomptarget_plugin:
LLVM/OpenMP Target Host Runtime Plugins (``libomptarget.rtl.XXXX``)
More information about the Openmp-commits
mailing list