[Offload] Implement double free (and other allocation error) reporting #100261
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
llvmbot
added
openmp:libomp
|
@llvm/pr-subscribers-offload Author: Johannes Doerfert (jdoerfert) ChangesAs a first step towards a GPU sanitizer we now can track allocations and deallocations in order to report double frees, and other problems during deallocation. Patch is 23.07 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/100261.diff 12 Files Affected:
diff --git a/offload/cmake/OpenMPTesting.cmake b/offload/cmake/OpenMPTesting.cmake
index 11eafeb764260..2ba0c06c264b8 100644
--- a/offload/cmake/OpenMPTesting.cmake
+++ b/offload/cmake/OpenMPTesting.cmake
@@ -47,6 +47,17 @@ function(find_standalone_test_dependencies)
set(ENABLE_CHECK_TARGETS FALSE PARENT_SCOPE)
return()
endif()
+
+ find_program(OFFLOAD_SYMBOLIZER_EXECUTABLE
+ NAMES llvm-symbolizer
+ PATHS ${OPENMP_LLVM_TOOLS_DIR})
+ if (NOT OFFLOAD_SYMBOLIZER_EXECUTABLE)
+ message(STATUS "Cannot find 'llvm-symbolizer'.")
+ message(STATUS "Please put 'llvm-symbolizer' in your PATH, set OFFLOAD_SYMBOLIZER_EXECUTABLE to its full path, or point OPENMP_LLVM_TOOLS_DIR to its directory.")
+ message(WARNING "The check targets will not be available!")
+ set(ENABLE_CHECK_TARGETS FALSE PARENT_SCOPE)
+ return()
+ endif()
endfunction()
if (${OPENMP_STANDALONE_BUILD})
@@ -71,6 +82,7 @@ else()
set(OPENMP_FILECHECK_EXECUTABLE ${LLVM_RUNTIME_OUTPUT_INTDIR}/FileCheck)
endif()
set(OPENMP_NOT_EXECUTABLE ${LLVM_RUNTIME_OUTPUT_INTDIR}/not)
+ set(OFFLOAD_SYMBOLIZER_EXECUTABLE ${LLVM_RUNTIME_OUTPUT_INTDIR}/llvm-symbolizer)
endif()
# Macro to extract information about compiler from file. (no own scope)
diff --git a/offload/plugins-nextgen/common/include/ErrorReporting.h b/offload/plugins-nextgen/common/include/ErrorReporting.h
new file mode 100644
index 0000000000000..53fb25effa5bc
--- /dev/null
+++ b/offload/plugins-nextgen/common/include/ErrorReporting.h
@@ -0,0 +1,200 @@
+//===- ErrorReporting.h - Helper to provide nice error messages ----- c++ -===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef OPENMP_LIBOMPTARGET_PLUGINS_NEXTGEN_COMMON_ERROR_REPORTING_H
+#define OPENMP_LIBOMPTARGET_PLUGINS_NEXTGEN_COMMON_ERROR_REPORTING_H
+
+#include "PluginInterface.h"
+
+#include "llvm/ADT/SmallString.h"
+#include "llvm/ADT/StringRef.h"
+#include "llvm/Support/ErrorHandling.h"
+
+#include <cstdio>
+#include <cstdlib>
+#include <string>
+
+namespace llvm {
+namespace omp {
+namespace target {
+namespace plugin {
+
+class ErrorReporter {
+ /// The banner printed at the beginning of an error report.
+ static constexpr auto ErrorBanner = "OFFLOAD ERROR: ";
+
+ /// Terminal color codes
+ ///
+ /// TODO: determine if the terminal supports colors.
+ ///@{
+ static constexpr auto Green = []() { return "\033[1m\033[32m"; };
+ static constexpr auto Blue = []() { return "\033[1m\033[34m"; };
+ static constexpr auto Red = []() { return "\033[1m\033[31m"; };
+ static constexpr auto Magenta = []() { return "\033[1m\033[35m"; };
+ static constexpr auto Cyan = []() { return "\033[1m\033[36m"; };
+ static constexpr auto Default = []() { return "\033[1m\033[0m"; };
+ ///@}
+
+ /// The size of the getBuffer() buffer.
+ static constexpr unsigned BufferSize = 1024;
+
+ /// Return a buffer of size BufferSize that can be used for formatting.
+ static char *getBuffer() {
+ static char *Buffer = nullptr;
+ if (!Buffer)
+ Buffer = reinterpret_cast<char *>(malloc(BufferSize));
+ return Buffer;
+ }
+
+ /// Return the device id as string, or n/a if not available.
+ static std::string getDeviceIdStr(GenericDeviceTy *Device) {
+ return Device ? std::to_string(Device->getDeviceId()) : "n/a";
+ }
+
+ /// Return a nice name for an TargetAllocTy.
+ static std::string getAllocTyName(TargetAllocTy Kind) {
+ switch (Kind) {
+ case TARGET_ALLOC_DEVICE_NON_BLOCKING:
+ case TARGET_ALLOC_DEFAULT:
+ case TARGET_ALLOC_DEVICE:
+ return "device memory";
+ case TARGET_ALLOC_HOST:
+ return "pinned host memory";
+ case TARGET_ALLOC_SHARED:
+ return "managed memory";
+ break;
+ }
+ llvm_unreachable("Unknown target alloc kind");
+ }
+
+ /// Return a C string after \p Format has been instantiated with \p Args.
+ template <typename... ArgsTy>
+ static const char *getCString(const char *Format, ArgsTy &&...Args) {
+ std::snprintf(getBuffer(), BufferSize, Format,
+ std::forward<ArgsTy>(Args)...);
+ return getBuffer();
+ }
+
+ /// Print \p Format, instantiated with \p Args to stderr.
+ /// TODO: Allow redirection into a file stream.
+ template <typename... ArgsTy>
+ static void print(const char *Format, ArgsTy &&...Args) {
+ fprintf(stderr, Format, std::forward<ArgsTy>(Args)...);
+ }
+
+ /// Pretty print a stack trace.
+ static void reportStackTrace(StringRef StackTrace) {
+ if (StackTrace.empty())
+ return;
+
+ SmallVector<StringRef> Lines, Parts;
+ StackTrace.split(Lines, "\n", /*MaxSplit=*/-1, /*KeepEmpty=*/false);
+ int Start = Lines.empty() || !Lines[0].contains("PrintStackTrace") ? 0 : 1;
+ for (int I = Start, E = Lines.size(); I < E; ++I) {
+ auto Line = Lines[I];
+ Parts.clear();
+ Line.split(Parts, " ", /*MaxSplit=*/2);
+ if (Parts.size() != 3 || Parts[0].size() < 2 || Parts[0][0] != '#') {
+ print("%s\n", Line.str().c_str());
+ continue;
+ }
+ unsigned FrameIdx = std::stoi(Parts[0].drop_front(1).str());
+ if (Start)
+ FrameIdx -= 1;
+ print(" %s%s%s%u %s%s%s %s\n", Magenta(),
+ Parts[0].take_front().str().c_str(), Green(), FrameIdx, Blue(),
+ Parts[1].str().c_str(), Default(), Parts[2].str().c_str());
+ }
+
+ printf("\n");
+ }
+
+ /// Report an error.
+ static void reportError(const char *Message, StringRef StackTrace) {
+ print("%s%s%s\n%s", Red(), ErrorBanner, Message, Default());
+ reportStackTrace(StackTrace);
+ }
+
+ /// Report information about an allocation associated with \p ATI.
+ static void reportAllocationInfo(AllocationTraceInfoTy *ATI) {
+ if (!ATI)
+ return;
+
+ if (!ATI->DeallocationTrace.empty()) {
+ print("%s%s%s\n%s", Cyan(), "Last deallocation:", Default());
+ reportStackTrace(ATI->DeallocationTrace);
+ }
+
+ if (ATI->HostPtr)
+ print("%sLast allocation of size %lu for host pointer %p:\n%s", Cyan(),
+ ATI->Size, ATI->HostPtr, Default());
+ else
+ print("%sLast allocation of size %lu:\n%s", Cyan(), ATI->Size, Default());
+ reportStackTrace(ATI->AllocationTrace);
+ if (!ATI->LastAllocationInfo)
+ return;
+
+ unsigned I = 0;
+ print("%sPrior allocations with the same base pointer:", Cyan());
+ while (ATI->LastAllocationInfo) {
+ print("\n%s", Default());
+ ATI = ATI->LastAllocationInfo;
+ print("%s #%u Prior deallocation of size %lu:\n%s", Cyan(), I, ATI->Size,
+ Default());
+ reportStackTrace(ATI->DeallocationTrace);
+ if (ATI->HostPtr)
+ print("%s #%u Prior allocation for host pointer %p:\n%s", Cyan(), I,
+ ATI->HostPtr, Default());
+ else
+ print("%s #%u Prior allocation:\n%s", Cyan(), I, Default());
+ reportStackTrace(ATI->AllocationTrace);
+ ++I;
+ }
+ }
+
+public:
+ /// Check if the deallocation of \p DevicePtr is valid given \p ATI. Stores \p
+ /// StackTrace to \p ATI->DeallocationTrace if there was no error.
+ static void checkDeallocation(GenericDeviceTy *Device, void *DevicePtr,
+ TargetAllocTy Kind, AllocationTraceInfoTy *ATI,
+ std::string &StackTrace) {
+#define DEALLOCATION_ERROR(Format, ...) \
+ reportError(getCString(Format, __VA_ARGS__), StackTrace); \
+ reportAllocationInfo(ATI); \
+ abort();
+
+ if (!ATI) {
+ DEALLOCATION_ERROR("deallocation of non-allocated %s: %p",
+ getAllocTyName(Kind).c_str(), DevicePtr);
+ }
+
+ if (!ATI->DeallocationTrace.empty()) {
+ DEALLOCATION_ERROR("double-free of %s: %p", getAllocTyName(Kind).c_str(),
+ DevicePtr);
+ }
+
+ if (ATI->Kind != Kind) {
+ DEALLOCATION_ERROR("deallocation requires %s but allocation was %s: %p",
+ getAllocTyName(Kind).c_str(),
+ getAllocTyName(ATI->Kind).c_str(), DevicePtr);
+ }
+
+ ATI->DeallocationTrace = StackTrace;
+
+#undef DEALLOCATION_ERROR
+ }
+};
+
+} // namespace plugin
+} // namespace target
+} // namespace omp
+} // namespace llvm
+
+#endif // OPENMP_LIBOMPTARGET_PLUGINS_NEXTGEN_COMMON_ERROR_REPORTING_H
diff --git a/offload/plugins-nextgen/common/include/PluginInterface.h b/offload/plugins-nextgen/common/include/PluginInterface.h
index 973add0ba1000..3ceb28e4d3ad3 100644
--- a/offload/plugins-nextgen/common/include/PluginInterface.h
+++ b/offload/plugins-nextgen/common/include/PluginInterface.h
@@ -19,6 +19,7 @@
#include <shared_mutex>
#include <vector>
+#include "ExclusiveAccess.h"
#include "Shared/APITypes.h"
#include "Shared/Debug.h"
#include "Shared/Environment.h"
@@ -382,6 +383,32 @@ struct GenericKernelTy {
bool IsBareKernel = false;
};
+/// Information about an allocation, when it has been allocated, and when/if it
+/// has been deallocated, for error reporting purposes.
+struct AllocationTraceInfoTy {
+
+ /// The stack trace of the allocation itself.
+ std::string AllocationTrace;
+
+ /// The stack trace of the deallocation, or empty.
+ std::string DeallocationTrace;
+
+ /// The allocated device pointer.
+ void *DevicePtr = nullptr;
+
+ /// The corresponding host pointer (can be null).
+ void *HostPtr = nullptr;
+
+ /// The size of the allocation.
+ uint64_t Size = 0;
+
+ /// The kind of the allocation.
+ TargetAllocTy Kind = TargetAllocTy::TARGET_ALLOC_DEFAULT;
+
+ /// Information about the last allocation at this address, if any.
+ AllocationTraceInfoTy *LastAllocationInfo = nullptr;
+};
+
/// Class representing a map of host pinned allocations. We track these pinned
/// allocations, so memory tranfers invloving these buffers can be optimized.
class PinnedAllocationMapTy {
@@ -866,6 +893,10 @@ struct GenericDeviceTy : public DeviceAllocatorTy {
/// Reference to the underlying plugin that created this device.
GenericPluginTy &Plugin;
+ /// Map to record when allocations have been performed, and when they have
+ /// been deallocated, both for error reporting purposes.
+ ProtectedObj<DenseMap<void *, AllocationTraceInfoTy *>> AllocationTraces;
+
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
@@ -916,6 +947,11 @@ struct GenericDeviceTy : public DeviceAllocatorTy {
UInt32Envar OMPX_InitialNumStreams;
UInt32Envar OMPX_InitialNumEvents;
+ /// Environment variable to determine if stack traces for allocations and
+ /// deallocations are tracked.
+ BoolEnvar OMPX_TrackAllocationTraces =
+ BoolEnvar("OFFLOAD_TRACK_ALLOCATION_TRACES", false);
+
/// Array of images loaded into the device. Images are automatically
/// deallocated by the allocator.
llvm::SmallVector<DeviceImageTy *> LoadedImages;
diff --git a/offload/plugins-nextgen/common/src/PluginInterface.cpp b/offload/plugins-nextgen/common/src/PluginInterface.cpp
index 118265973f327..f9b8df0f6bf21 100644
--- a/offload/plugins-nextgen/common/src/PluginInterface.cpp
+++ b/offload/plugins-nextgen/common/src/PluginInterface.cpp
@@ -14,6 +14,7 @@
#include "Shared/Debug.h"
#include "Shared/Environment.h"
+#include "ErrorReporting.h"
#include "GlobalHandler.h"
#include "JIT.h"
#include "Utils/ELF.h"
@@ -30,6 +31,8 @@
#include "llvm/Support/JSON.h"
#include "llvm/Support/MathExtras.h"
#include "llvm/Support/MemoryBuffer.h"
+#include "llvm/Support/Signals.h"
+#include "llvm/Support/raw_ostream.h"
#include <cstdint>
#include <limits>
@@ -1337,6 +1340,25 @@ Expected<void *> GenericDeviceTy::dataAlloc(int64_t Size, void *HostPtr,
if (auto Err = PinnedAllocs.registerHostBuffer(Alloc, Alloc, Size))
return std::move(Err);
+ // Keep track of the allocation stack if we track allocation traces.
+ if (OMPX_TrackAllocationTraces) {
+ std::string StackTrace;
+ llvm::raw_string_ostream OS(StackTrace);
+ llvm::sys::PrintStackTrace(OS);
+
+ AllocationTraceInfoTy *ATI = new AllocationTraceInfoTy();
+ ATI->AllocationTrace = std::move(StackTrace);
+ ATI->DevicePtr = Alloc;
+ ATI->HostPtr = HostPtr;
+ ATI->Size = Size;
+ ATI->Kind = Kind;
+
+ auto AllocationTraceMap = AllocationTraces.getExclusiveAccessor();
+ auto *&MapATI = (*AllocationTraceMap)[Alloc];
+ ATI->LastAllocationInfo = MapATI;
+ MapATI = ATI;
+ }
+
return Alloc;
}
@@ -1345,6 +1367,21 @@ Error GenericDeviceTy::dataDelete(void *TgtPtr, TargetAllocTy Kind) {
if (Plugin.getRecordReplay().isRecordingOrReplaying())
return Plugin::success();
+ // Keep track of the deallocation stack if we track allocation traces.
+ if (OMPX_TrackAllocationTraces) {
+ AllocationTraceInfoTy *ATI = nullptr;
+ {
+ auto AllocationTraceMap = AllocationTraces.getExclusiveAccessor();
+ ATI = (*AllocationTraceMap)[TgtPtr];
+ }
+
+ std::string StackTrace;
+ llvm::raw_string_ostream OS(StackTrace);
+ llvm::sys::PrintStackTrace(OS);
+
+ ErrorReporter::checkDeallocation(this, TgtPtr, Kind, ATI, StackTrace);
+ }
+
int Res;
switch (Kind) {
case TARGET_ALLOC_DEFAULT:
diff --git a/offload/src/omptarget.cpp b/offload/src/omptarget.cpp
index 9bca8529c5ee3..3b627d257a069 100644
--- a/offload/src/omptarget.cpp
+++ b/offload/src/omptarget.cpp
@@ -462,7 +462,9 @@ void targetFreeExplicit(void *DevicePtr, int DeviceNum, int Kind,
FATAL_MESSAGE(DeviceNum, "%s", toString(DeviceOrErr.takeError()).c_str());
if (DeviceOrErr->deleteData(DevicePtr, Kind) == OFFLOAD_FAIL)
- FATAL_MESSAGE(DeviceNum, "%s", "Failed to deallocate device ptr");
+ FATAL_MESSAGE(DeviceNum, "%s",
+ "Failed to deallocate device ptr. Set "
+ "OFFLOAD_TRACK_ALLOCATION_TRACES=1 to track allocations.");
DP("omp_target_free deallocated device ptr\n");
}
diff --git a/offload/test/lit.cfg b/offload/test/lit.cfg
index 99d7defdb9e11..eb729ef40418a 100644
--- a/offload/test/lit.cfg
+++ b/offload/test/lit.cfg
@@ -415,3 +415,4 @@ config.substitutions.append(("%flags_clang", config.test_flags_clang))
config.substitutions.append(("%flags_flang", config.test_flags_flang))
config.substitutions.append(("%flags", config.test_flags))
config.substitutions.append(("%not", config.libomptarget_not))
+config.substitutions.append(("%symbolizer", config.offload_symbolizer))
diff --git a/offload/test/lit.site.cfg.in b/offload/test/lit.site.cfg.in
index 43751970cac27..66415972feff4 100644
--- a/offload/test/lit.site.cfg.in
+++ b/offload/test/lit.site.cfg.in
@@ -23,6 +23,7 @@ config.libomptarget_all_targets = "@LIBOMPTARGET_ALL_TARGETS@".split()
config.libomptarget_current_target = "@CURRENT_TARGET@"
config.libomptarget_filecheck = "@OPENMP_FILECHECK_EXECUTABLE@"
config.libomptarget_not = "@OPENMP_NOT_EXECUTABLE@"
+config.offload_symbolizer = "@OFFLOAD_SYMBOLIZER_EXECUTABLE@"
config.libomptarget_debug = @LIBOMPTARGET_DEBUG@
config.has_libomptarget_ompt = @LIBOMPTARGET_OMPT_SUPPORT@
config.libomptarget_has_libc = @LIBOMPTARGET_GPU_LIBC_SUPPORT@
diff --git a/offload/test/sanitizer/double_free.c b/offload/test/sanitizer/double_free.c
new file mode 100644
index 0000000000000..e6dd94d2e09be
--- /dev/null
+++ b/offload/test/sanitizer/double_free.c
@@ -0,0 +1,68 @@
+// clang-format off
+// RUN: %libomptarget-compileopt-generic
+// RUN: %not --crash env -u LLVM_DISABLE_SYMBOLIZATION LLVM_SYMBOLIZER_PATH=%symbolizer OFFLOAD_TRACK_ALLOCATION_TRACES=1 %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK,NDEBG
+// RUN: %libomptarget-compileopt-generic -g
+// RUN: %not --crash env -u LLVM_DISABLE_SYMBOLIZATION LLVM_SYMBOLIZER_PATH=%symbolizer OFFLOAD_TRACK_ALLOCATION_TRACES=1 %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK,DEBUG
+// clang-format on
+
+// 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) {
+ void *Ptr1 = omp_target_alloc(8, 0);
+ omp_target_free(Ptr1, 0);
+ void *Ptr2 = omp_target_alloc(8, 0);
+ omp_target_free(Ptr2, 0);
+ void *Ptr3 = omp_target_alloc(8, 0);
+ omp_target_free(Ptr3, 0);
+ omp_target_free(Ptr2, 0);
+}
+
+// CHECK: OFFLOAD ERROR: double-free of device memory: 0x
+// CHECK: {{.*}}dataDelete
+// CHECK: omp_target_free
+// NDEBG: main
+// DEBUG: main {{.*}}double_free.c:25
+//
+// CHECK: Last deallocation:
+// CHECK: {{.*}}dataDelete
+// CHECK: omp_target_free
+// NDEBG: main
+// DEBUG: main {{.*}}double_free.c:24
+//
+// CHECK: Last allocation of size 8:
+// CHECK: {{.*}}dataAlloc
+// CHECK: omp_target_alloc
+// NDEBG: main
+// DEBUG: main {{.*}}double_free.c:23
+//
+// CHECK: Prior allocations with the same base pointer:
+// CHECK: #0 Prior deallocation of size 8:
+// CHECK: {{.*}}dataDelete
+// CHECK: omp_target_free
+// NDEBG: main
+// DEBUG: main {{.*}}double_free.c:22
+//
+// CHECK: #0 Prior allocation:
+// CHECK: {{.*}}dataAlloc
+// CHECK: omp_target_alloc
+// NDEBG: main
+// DEBUG: main {{.*}}double_free.c:20
+//
+// CHECK: #1 Prior deallocation of size 8:
+// CHECK: {{.*}}dataDelete
+// CHECK: omp_target_free
+// NDEBG: main
+// DEBUG: main {{.*}}double_free.c:20
+//
+// CHECK: #1 Prior allocation:
+// CHECK: {{.*}}dataAlloc
+// CHECK: omp_target_alloc
+// NDEBG: main
+// DEBUG: main {{.*}}double_free.c:19
diff --git a/offload/test/sanitizer/free_host_ptr.c b/offload/test/sanitizer/free_host_ptr.c
new file mode 100644
index 0000000000000..65a7bfa97dca0
--- /dev/null
+++ b/offload/test/sanitizer/free_host_ptr.c
@@ -0,0 +1,25 @@
+// clang-format off
+// RUN: %libomptarget-compileopt-generic
+// RUN: %not --crash env -u LLVM_DISABLE_SYMBOLIZATION LLVM_SYMBOLIZER_PATH=%symbolizer OFFLOAD_TRACK_ALLOCATION_TRACES=1 %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK,NDEBG
+// RUN: %libomptarget-compileopt-generic -g
+// RUN: %not --crash env -u LLVM_DISABLE_SYMBOLIZATION LLVM_SYMBOLIZER_PATH=%symbolizer OFFLOAD_TRACK_ALLOCATION_TRACES=1 %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK,DEBUG
+// clang-format on
+
+// 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 X;
+ omp_target_free(&X, 0);
+}
+
+// CHECK: OFFLOAD ERROR: deallocation of non-allocated device memory: 0x
+// CHECK: {{.*}}dataDelete
+// NDEBG: {{.*}}main
+// DEBUG: {{.*}}main {{.*}}free_host_ptr.c:20
diff --git a/offload/test/sanitizer/free_wrong_ptr_kind.c b/offload/test/sanitizer/free_wrong_ptr_kind.c
new file mode 100644
index 0000000000000..040a269956643
--- /dev/null
+++ b/offload/test/sanitizer/free_wrong_ptr_kind.c
@@ -0,0 +1,35 @@
+// clang-format off
+// RUN: %libomptarget-compileopt-generic
+// RUN: %not --crash env -u LLVM_DISABLE_SYMBOLIZATION LLVM_SYMBOLIZER_PATH=%symbolizer OFFLOAD_TRACK_ALLOCATION_TRACES=1 %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK,NDEBG
+// RUN: %libomptarget-compileopt-generic -g
+// RUN: %not --crash env -u LLVM_DISABLE_SYMBOLIZATION LLVM_SYMBOLIZER_PATH=%symbolizer OFFLOAD_TRACK_ALLOCATION_TRACES=1 %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK,DEBUG
+// clang-format on
+
+// 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>
+
+void *llvm_omp_target_alloc_host(size_t Size, int DeviceNum);
+
+int main(void) {
+ void *P = llvm_omp_target_alloc_host(8, 0);
+ omp_target_free(P, 0);
+}
+
+// clang-format off
+// CHECK: OFFLOAD ERROR: deallocation requires...
[truncated]
|
jdoerfert
force-pushed
the
pr/double_free_reporting
branch
from
17e3632 to
c3a3bab
Compare
shiltian
reviewed
Jul 23, 2024
shiltian
reviewed
Jul 23, 2024
jdoerfert
force-pushed
the
pr/double_free_reporting
branch
from
0d2b626 to
e44480e
Compare
shiltian
reviewed
Jul 25, 2024
jdoerfert
force-pushed
the
pr/double_free_reporting
branch
from
e44480e to
388df39
Compare
jhuber6
reviewed
Jul 25, 2024
shiltian
reviewed
Jul 25, 2024
jdoerfert
force-pushed
the
pr/double_free_reporting
branch
from
388df39 to
17e6fd7
Compare
shiltian
reviewed
Jul 25, 2024
jdoerfert
force-pushed
the
pr/double_free_reporting
branch
from
17e6fd7 to
f5bf4f2
Compare
jhuber6
reviewed
Jul 25, 2024
| } | ||
|
|
||
| /// Pretty print a stack trace. | ||
| static void reportStackTrace(StringRef StackTrace) { |
There was a problem hiding this comment.
I thought that LLVM had their own format for presenting stack traces. I guess we want to add more color and information than what it exports?
There was a problem hiding this comment.
We use LLVM to get the stack trace as string.
From there, we 1) remove the first line (nobody want's to know we are printing the stack trace) 2) add color and 3) decide where to print it to (for now stderr).
Top is LLVMs stack trace as printed to errs(), bottom what we print with this.
jplehr
reviewed
Jul 26, 2024
|
|
||
| int main(void) { | ||
| void *Ptr1 = omp_target_alloc(8, 0); | ||
| #pragma omp parallel |
There was a problem hiding this comment.
Does something in lit limit this to a certain number of host threads? I wonder how likely it is that we flood a buildbot in case we have several of such tests, which may run in parallel.
There was a problem hiding this comment.
I don't think lit limits it, no. I can, and feel free to limit other tests if there is a problem. We had a student work on a way for lit to understand parallelism and issue them appropriately, but it never went in (and I don't know where the code is).
As a first step towards a GPU sanitizer we now can track allocations and deallocations in order to report double frees, and other problems during deallocation.
jdoerfert
force-pushed
the
pr/double_free_reporting
branch
from
f5bf4f2 to
3b28b65
Compare
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.
As a first step towards a GPU sanitizer we now can track allocations and deallocations in order to report double frees, and other problems during deallocation.