[llvm] [OpenMP][Offload] Continue to update libomptarget debug messages (PR #170425)
Alex Duran via llvm-commits
llvm-commits at lists.llvm.org
Wed Dec 10 05:13:37 PST 2025
https://github.com/adurang updated https://github.com/llvm/llvm-project/pull/170425
>From f63c8e6027290d797a99a00c6574f08d68cb09b1 Mon Sep 17 00:00:00 2001
From: Alex Duran <alejandro.duran at intel.com>
Date: Tue, 2 Dec 2025 09:56:37 +0100
Subject: [PATCH 1/7] [OFFLOAD][LIBOMPTARGET] Add compatibility support; start
to update messages
---
offload/include/Shared/Debug.h | 327 +++++++++++++++++-----------
offload/libomptarget/OffloadRTL.cpp | 8 +-
offload/libomptarget/device.cpp | 14 +-
3 files changed, 207 insertions(+), 142 deletions(-)
diff --git a/offload/include/Shared/Debug.h b/offload/include/Shared/Debug.h
index 41613a37c3548..feba3c9dddb29 100644
--- a/offload/include/Shared/Debug.h
+++ b/offload/include/Shared/Debug.h
@@ -39,6 +39,7 @@
#define OMPTARGET_SHARED_DEBUG_H
#include <atomic>
+#include <cstdarg>
#include <mutex>
#include <string>
@@ -78,17 +79,6 @@ inline std::atomic<uint32_t> &getInfoLevelInternal() {
inline uint32_t getInfoLevel() { return getInfoLevelInternal().load(); }
-inline uint32_t getDebugLevel() {
- static uint32_t DebugLevel = 0;
- static std::once_flag Flag{};
- std::call_once(Flag, []() {
- if (char *EnvStr = getenv("LIBOMPTARGET_DEBUG"))
- DebugLevel = std::stoi(EnvStr);
- });
-
- return DebugLevel;
-}
-
#undef USED
#undef GCC_VERSION
@@ -147,46 +137,11 @@ inline uint32_t getDebugLevel() {
fprintf(_stdDst, __VA_ARGS__); \
} while (0)
-// Debugging messages
-#ifdef OMPTARGET_DEBUG
-#include <stdio.h>
-
-#define DEBUGP(prefix, ...) \
- { \
- fprintf(stderr, "%s --> ", prefix); \
- fprintf(stderr, __VA_ARGS__); \
- }
-
-/// Emit a message for debugging
-#define DP(...) \
- do { \
- if (getDebugLevel() > 0) { \
- DEBUGP(DEBUG_PREFIX, __VA_ARGS__); \
- } \
- } while (false)
-
-/// Emit a message for debugging or failure if debugging is disabled
-#define REPORT(...) \
- do { \
- if (getDebugLevel() > 0) { \
- DP(__VA_ARGS__); \
- } else { \
- FAILURE_MESSAGE(__VA_ARGS__); \
- } \
- } while (false)
-#else
-#define DEBUGP(prefix, ...) \
- {}
-#define DP(...) \
- {}
-#define REPORT(...) FAILURE_MESSAGE(__VA_ARGS__);
-#endif // OMPTARGET_DEBUG
-
/// Emit a message giving the user extra information about the runtime if
#define INFO(_flags, _id, ...) \
do { \
- if (getDebugLevel() > 0) { \
- DEBUGP(DEBUG_PREFIX, __VA_ARGS__); \
+ if (::llvm::offload::debug::isDebugEnabled()) { \
+ DP(__VA_ARGS__); \
} else if (getInfoLevel() & _flags) { \
INFO_MESSAGE(_id, __VA_ARGS__); \
} \
@@ -203,17 +158,92 @@ inline uint32_t getDebugLevel() {
namespace llvm::offload::debug {
-#ifdef OMPTARGET_DEBUG
+/// A raw_ostream that tracks `\n` and print the prefix after each
+/// newline. Based on raw_ldbg_ostream from Support/DebugLog.h
+class LLVM_ABI odbg_ostream final : public raw_ostream {
+public:
+ enum IfLevel : uint32_t;
+ enum OnlyLevel : uint32_t;
-struct DebugFilter {
- StringRef Type;
- uint32_t Level;
-};
+private:
+ std::string Prefix;
+ raw_ostream &Os;
+ uint32_t BaseLevel;
+ bool ShouldPrefixNextString;
+ bool ShouldEmitNewLineOnDestruction;
+ bool NeedEndNewLine = false;
-struct DebugSettings {
- bool Enabled = false;
- uint32_t DefaultLevel = 1;
- llvm::SmallVector<DebugFilter> Filters;
+ /// If the stream is muted, writes to it are ignored
+ bool Muted = false;
+
+ /// Split the line on newlines and insert the prefix before each
+ /// newline. Forward everything to the underlying stream.
+ void write_impl(const char *Ptr, size_t Size) final {
+ if (Muted)
+ return;
+
+ NeedEndNewLine = false;
+ auto Str = StringRef(Ptr, Size);
+ auto Eol = Str.find('\n');
+ // Handle `\n` occurring in the string, ensure to print the prefix at the
+ // beginning of each line.
+ while (Eol != StringRef::npos) {
+ // Take the line up to the newline (including the newline).
+ StringRef Line = Str.take_front(Eol + 1);
+ if (!Line.empty())
+ writeWithPrefix(Line);
+ // We printed a newline, record here to print a prefix.
+ ShouldPrefixNextString = true;
+ Str = Str.drop_front(Eol + 1);
+ Eol = Str.find('\n');
+ }
+ if (!Str.empty()) {
+ writeWithPrefix(Str);
+ NeedEndNewLine = true;
+ }
+ }
+ void emitPrefix() { Os.write(Prefix.c_str(), Prefix.size()); }
+ void writeWithPrefix(StringRef Str) {
+ if (ShouldPrefixNextString) {
+ emitPrefix();
+ ShouldPrefixNextString = false;
+ }
+ Os.write(Str.data(), Str.size());
+ }
+
+public:
+ explicit odbg_ostream(std::string Prefix, raw_ostream &Os, uint32_t BaseLevel,
+ bool ShouldPrefixNextString = true,
+ bool ShouldEmitNewLineOnDestruction = true)
+ : Prefix(std::move(Prefix)), Os(Os), BaseLevel(BaseLevel),
+ ShouldPrefixNextString(ShouldPrefixNextString),
+ ShouldEmitNewLineOnDestruction(ShouldEmitNewLineOnDestruction) {
+ SetUnbuffered();
+ }
+ ~odbg_ostream() final {
+ if (ShouldEmitNewLineOnDestruction && NeedEndNewLine)
+ Os << '\n';
+ }
+ odbg_ostream(const odbg_ostream &) = delete;
+ odbg_ostream &operator=(const odbg_ostream &) = delete;
+ odbg_ostream(odbg_ostream &&other) : Os(other.Os) {
+ Prefix = std::move(other.Prefix);
+ BaseLevel = other.BaseLevel;
+ ShouldPrefixNextString = other.ShouldPrefixNextString;
+ ShouldEmitNewLineOnDestruction = other.ShouldEmitNewLineOnDestruction;
+ NeedEndNewLine = other.NeedEndNewLine;
+ Muted = other.Muted;
+ }
+
+ /// Forward the current_pos method to the underlying stream.
+ uint64_t current_pos() const final { return Os.tell(); }
+
+ /// Some of the `<<` operators expect an lvalue, so we trick the type
+ /// system.
+ odbg_ostream &asLvalue() { return *this; }
+
+ void shouldMute(const IfLevel Filter) { Muted = Filter > BaseLevel; }
+ void shouldMute(const OnlyLevel Filter) { Muted = BaseLevel != Filter; }
};
/// dbgs - Return a circular-buffered debug stream.
@@ -228,6 +258,19 @@ struct DebugSettings {
return thestrm.strm;
}
+#ifdef OMPTARGET_DEBUG
+
+struct DebugFilter {
+ StringRef Type;
+ uint32_t Level;
+};
+
+struct DebugSettings {
+ bool Enabled = false;
+ uint32_t DefaultLevel = 1;
+ llvm::SmallVector<DebugFilter> Filters;
+};
+
[[maybe_unused]] static DebugFilter parseDebugFilter(StringRef Filter) {
size_t Pos = Filter.find(':');
if (Pos == StringRef::npos)
@@ -309,80 +352,6 @@ shouldPrintDebug(const char *Component, const char *Type, uint32_t &Level) {
return false;
}
-/// A raw_ostream that tracks `\n` and print the prefix after each
-/// newline. Based on raw_ldbg_ostream from Support/DebugLog.h
-class LLVM_ABI odbg_ostream final : public raw_ostream {
-public:
- enum IfLevel : uint32_t;
- enum OnlyLevel : uint32_t;
-
-private:
- std::string Prefix;
- raw_ostream &Os;
- uint32_t BaseLevel;
- bool ShouldPrefixNextString;
- bool ShouldEmitNewLineOnDestruction;
-
- /// If the stream is muted, writes to it are ignored
- bool Muted = false;
-
- /// Split the line on newlines and insert the prefix before each
- /// newline. Forward everything to the underlying stream.
- void write_impl(const char *Ptr, size_t Size) final {
- if (Muted)
- return;
-
- auto Str = StringRef(Ptr, Size);
- auto Eol = Str.find('\n');
- // Handle `\n` occurring in the string, ensure to print the prefix at the
- // beginning of each line.
- while (Eol != StringRef::npos) {
- // Take the line up to the newline (including the newline).
- StringRef Line = Str.take_front(Eol + 1);
- if (!Line.empty())
- writeWithPrefix(Line);
- // We printed a newline, record here to print a prefix.
- ShouldPrefixNextString = true;
- Str = Str.drop_front(Eol + 1);
- Eol = Str.find('\n');
- }
- if (!Str.empty())
- writeWithPrefix(Str);
- }
- void emitPrefix() { Os.write(Prefix.c_str(), Prefix.size()); }
- void writeWithPrefix(StringRef Str) {
- if (ShouldPrefixNextString) {
- emitPrefix();
- ShouldPrefixNextString = false;
- }
- Os.write(Str.data(), Str.size());
- }
-
-public:
- explicit odbg_ostream(std::string Prefix, raw_ostream &Os, uint32_t BaseLevel,
- bool ShouldPrefixNextString = true,
- bool ShouldEmitNewLineOnDestruction = false)
- : Prefix(std::move(Prefix)), Os(Os), BaseLevel(BaseLevel),
- ShouldPrefixNextString(ShouldPrefixNextString),
- ShouldEmitNewLineOnDestruction(ShouldEmitNewLineOnDestruction) {
- SetUnbuffered();
- }
- ~odbg_ostream() final {
- if (ShouldEmitNewLineOnDestruction)
- Os << '\n';
- }
-
- /// Forward the current_pos method to the underlying stream.
- uint64_t current_pos() const final { return Os.tell(); }
-
- /// Some of the `<<` operators expect an lvalue, so we trick the type
- /// system.
- odbg_ostream &asLvalue() { return *this; }
-
- void shouldMute(const IfLevel Filter) { Muted = Filter > BaseLevel; }
- void shouldMute(const OnlyLevel Filter) { Muted = BaseLevel != Filter; }
-};
-
/// Compute the prefix for the debug log in the form of:
/// "Component --> "
[[maybe_unused]] static std::string computePrefix(StringRef Component,
@@ -463,6 +432,8 @@ static inline raw_ostream &operator<<(raw_ostream &Os,
#else
+inline bool isDebugEnabled() { return false; }
+
#define ODBG_NULL \
for (bool _c = false; _c; _c = false) \
::llvm::nulls()
@@ -479,4 +450,98 @@ static inline raw_ostream &operator<<(raw_ostream &Os,
} // namespace llvm::offload::debug
+namespace llvm::omptarget::debug {
+using namespace llvm::offload::debug;
+
+enum OmpDebugLevel : uint32_t {
+ ODL_Default = 1,
+ ODL_Error = ODL_Default,
+ ODL_Detailed = 2,
+ ODL_Verbose = 3,
+ ODL_VeryVerbose = 4,
+ ODL_Dumping = 5
+};
+
+/* Debug types to use in libomptarget */
+constexpr const char *ODT_Init = "Init";
+constexpr const char *ODT_Mapping = "Mapping";
+constexpr const char *ODT_Kernel = "Kernel";
+constexpr const char *ODT_DataTransfer = "DataTransfer";
+constexpr const char *ODT_Sync = "Sync";
+constexpr const char *ODT_Deinit = "Deinit";
+constexpr const char *ODT_Error = "Error";
+constexpr const char *ODT_KernelArgs = "KernelArgs";
+constexpr const char *ODT_MappingExists = "MappingExists";
+constexpr const char *ODT_DumpTable = "DumpTable";
+constexpr const char *ODT_MappingChanged = "MappingChanged";
+constexpr const char *ODT_PluginKernel = "PluginKernel";
+constexpr const char *ODT_EmptyMapping = "EmptyMapping";
+
+static inline odbg_ostream reportErrorStream() {
+#ifdef OMPTARGET_DEBUG
+ if (::llvm::offload::debug::isDebugEnabled()) {
+ uint32_t RealLevel = ODL_Error;
+ if (::llvm::offload::debug::shouldPrintDebug(GETNAME(TARGET_NAME),
+ (ODT_Error), RealLevel))
+ return odbg_ostream{
+ ::llvm::offload::debug::computePrefix(DEBUG_PREFIX, ODT_Error),
+ ::llvm::offload::debug::dbgs(), RealLevel};
+ else
+ return odbg_ostream{"", ::llvm::nulls(), 1};
+ }
+#endif
+ return odbg_ostream{GETNAME(TARGET_NAME) " error: ",
+ ::llvm::offload::debug::dbgs(), ODL_Error};
+}
+
+#ifdef OMPTARGET_DEBUG
+// Deprecated debug print macros
+[[maybe_unused]] static std::string formatToStr(const char *format, ...) {
+ va_list args;
+ va_start(args, format);
+ size_t len = std::vsnprintf(NULL, 0, format, args);
+ va_end(args);
+ llvm::SmallVector<char, 128> vec(len + 1);
+ va_start(args, format);
+ std::vsnprintf(&vec[0], len + 1, format, args);
+ va_end(args);
+ return &vec[0];
+}
+
+// helper macro to support old DP and REPORT macros with printf syntax
+#define FORMAT_TO_STR(Format, ...) \
+ ::llvm::omptarget::debug::formatToStr(Format __VA_OPT__(, ) __VA_ARGS__)
+
+#define DP(...) ODBG() << FORMAT_TO_STR(__VA_ARGS__);
+
+#define REPORT_INT_OLD(...) \
+ do { \
+ if (::llvm::offload::debug::isDebugEnabled()) { \
+ ODBG(ODT_Error, ODL_Error) << FORMAT_TO_STR(__VA_ARGS__); \
+ } else { \
+ FAILURE_MESSAGE(__VA_ARGS__); \
+ } \
+ } while (false)
+
+#else
+#define DP(...) \
+ { \
+ }
+#define REPORT_INT_OLD(...) FAILURE_MESSAGE(__VA_ARGS__);
+#endif // OMPTARGET_DEBUG
+
+// This is used for the new style REPORT macro
+#define REPORT_INT() ::llvm::omptarget::debug::reportErrorStream()
+
+// Make REPORT compatible with old and new syntax
+#define REPORT(...) REPORT_INT##__VA_OPT__(_OLD)(__VA_ARGS__)
+
+} // namespace llvm::omptarget::debug
+
+using namespace llvm::omptarget::debug;
+
+static inline int getDebugLevel() {
+ return ::llvm::offload::debug::isDebugEnabled();
+}
+
#endif // OMPTARGET_SHARED_DEBUG_H
diff --git a/offload/libomptarget/OffloadRTL.cpp b/offload/libomptarget/OffloadRTL.cpp
index 0ae325bf496d9..77c5768b62168 100644
--- a/offload/libomptarget/OffloadRTL.cpp
+++ b/offload/libomptarget/OffloadRTL.cpp
@@ -35,7 +35,7 @@ void initRuntime() {
RefCount++;
if (RefCount == 1) {
- ODBG() << "Init offload library!";
+ ODBG(ODT_Init) << "Init offload library!";
#ifdef OMPT_SUPPORT
// Initialize OMPT first
llvm::omp::target::ompt::connectLibrary();
@@ -54,12 +54,12 @@ void deinitRuntime() {
assert(PM && "Runtime not initialized");
if (RefCount == 1) {
- DP("Deinit offload library!\n");
+ ODBG(ODT_Deinit) << "Deinit offload library!";
// RTL deinitialization has started
RTLAlive = false;
while (RTLOngoingSyncs > 0) {
- DP("Waiting for ongoing syncs to finish, count: %d\n",
- RTLOngoingSyncs.load());
+ ODBG(ODT_Sync) << "Waiting for ongoing syncs to finish, count:"
+ << RTLOngoingSyncs.load();
std::this_thread::sleep_for(std::chrono::milliseconds(100));
}
PM->deinit();
diff --git a/offload/libomptarget/device.cpp b/offload/libomptarget/device.cpp
index ee36fbed935a5..5637a77508039 100644
--- a/offload/libomptarget/device.cpp
+++ b/offload/libomptarget/device.cpp
@@ -48,7 +48,7 @@ int HostDataToTargetTy::addEventIfNecessary(DeviceTy &Device,
void *Event = getEvent();
bool NeedNewEvent = Event == nullptr;
if (NeedNewEvent && Device.createEvent(&Event) != OFFLOAD_SUCCESS) {
- REPORT("Failed to create event\n");
+ REPORT() << "Failed to create event";
return OFFLOAD_FAIL;
}
@@ -56,7 +56,7 @@ int HostDataToTargetTy::addEventIfNecessary(DeviceTy &Device,
// know if the target support event. But if a target doesn't,
// recordEvent should always return success.
if (Device.recordEvent(Event, AsyncInfo) != OFFLOAD_SUCCESS) {
- REPORT("Failed to set dependence on event " DPxMOD "\n", DPxPTR(Event));
+ REPORT() << "Failed to set dependence on event " << Event;
return OFFLOAD_FAIL;
}
@@ -278,21 +278,21 @@ int32_t DeviceTy::dataFence(AsyncInfoTy &AsyncInfo) {
}
int32_t DeviceTy::notifyDataMapped(void *HstPtr, int64_t Size) {
- DP("Notifying about new mapping: HstPtr=" DPxMOD ", Size=%" PRId64 "\n",
- DPxPTR(HstPtr), Size);
+ ODBG(ODT_Mapping) << "Notifying about new mapping: HstPtr=" << HstPtr
+ << ", Size=" << Size;
if (RTL->data_notify_mapped(RTLDeviceID, HstPtr, Size)) {
- REPORT("Notifying about data mapping failed.\n");
+ REPORT() << "Notifying about data mapping failed.";
return OFFLOAD_FAIL;
}
return OFFLOAD_SUCCESS;
}
int32_t DeviceTy::notifyDataUnmapped(void *HstPtr) {
- DP("Notifying about an unmapping: HstPtr=" DPxMOD "\n", DPxPTR(HstPtr));
+ ODBG(ODT_Mapping) << "Notifying about an unmapping: HstPtr=" << HstPtr;
if (RTL->data_notify_unmapped(RTLDeviceID, HstPtr)) {
- REPORT("Notifying about data unmapping failed.\n");
+ REPORT() << "Notifying about data unmapping failed.";
return OFFLOAD_FAIL;
}
return OFFLOAD_SUCCESS;
>From 1305a3eeafe10920f87e5140a43cce64af8aa9a3 Mon Sep 17 00:00:00 2001
From: Alex Duran <alejandro.duran at intel.com>
Date: Tue, 2 Dec 2025 10:05:01 +0100
Subject: [PATCH 2/7] minor cleanup
---
offload/include/Shared/Debug.h | 10 +++-------
offload/libomptarget/OffloadRTL.cpp | 1 +
offload/libomptarget/device.cpp | 1 +
3 files changed, 5 insertions(+), 7 deletions(-)
diff --git a/offload/include/Shared/Debug.h b/offload/include/Shared/Debug.h
index feba3c9dddb29..3a7687b24d8d4 100644
--- a/offload/include/Shared/Debug.h
+++ b/offload/include/Shared/Debug.h
@@ -517,7 +517,9 @@ static inline odbg_ostream reportErrorStream() {
#define REPORT_INT_OLD(...) \
do { \
if (::llvm::offload::debug::isDebugEnabled()) { \
- ODBG(ODT_Error, ODL_Error) << FORMAT_TO_STR(__VA_ARGS__); \
+ ODBG(::llvm::omptarget::debug::ODT_Error, \
+ ::llvm::omptarget::debug::ODL_Error) \
+ << FORMAT_TO_STR(__VA_ARGS__); \
} else { \
FAILURE_MESSAGE(__VA_ARGS__); \
} \
@@ -538,10 +540,4 @@ static inline odbg_ostream reportErrorStream() {
} // namespace llvm::omptarget::debug
-using namespace llvm::omptarget::debug;
-
-static inline int getDebugLevel() {
- return ::llvm::offload::debug::isDebugEnabled();
-}
-
#endif // OMPTARGET_SHARED_DEBUG_H
diff --git a/offload/libomptarget/OffloadRTL.cpp b/offload/libomptarget/OffloadRTL.cpp
index 77c5768b62168..3dc37db1e1d67 100644
--- a/offload/libomptarget/OffloadRTL.cpp
+++ b/offload/libomptarget/OffloadRTL.cpp
@@ -19,6 +19,7 @@
#ifdef OMPT_SUPPORT
extern void llvm::omp::target::ompt::connectLibrary();
#endif
+using namespace llvm::omptarget::debug;
static std::mutex PluginMtx;
static uint32_t RefCount = 0;
diff --git a/offload/libomptarget/device.cpp b/offload/libomptarget/device.cpp
index 5637a77508039..055b901372a37 100644
--- a/offload/libomptarget/device.cpp
+++ b/offload/libomptarget/device.cpp
@@ -38,6 +38,7 @@ using namespace llvm::omp::target::ompt;
#endif
using namespace llvm::omp::target::plugin;
+using namespace llvm::omptarget::debug;
int HostDataToTargetTy::addEventIfNecessary(DeviceTy &Device,
AsyncInfoTy &AsyncInfo) const {
>From 845c5b2587a9ce358104aba2fccfcda5319bad17 Mon Sep 17 00:00:00 2001
From: Alex Duran <alejandro.duran at intel.com>
Date: Tue, 2 Dec 2025 10:07:44 +0100
Subject: [PATCH 3/7] update namespace name
---
offload/include/Shared/Debug.h | 12 ++++++------
offload/libomptarget/OffloadRTL.cpp | 2 +-
offload/libomptarget/device.cpp | 2 +-
3 files changed, 8 insertions(+), 8 deletions(-)
diff --git a/offload/include/Shared/Debug.h b/offload/include/Shared/Debug.h
index 3a7687b24d8d4..9e657e64484c0 100644
--- a/offload/include/Shared/Debug.h
+++ b/offload/include/Shared/Debug.h
@@ -450,7 +450,7 @@ inline bool isDebugEnabled() { return false; }
} // namespace llvm::offload::debug
-namespace llvm::omptarget::debug {
+namespace llvm::omp::target::debug {
using namespace llvm::offload::debug;
enum OmpDebugLevel : uint32_t {
@@ -510,15 +510,15 @@ static inline odbg_ostream reportErrorStream() {
// helper macro to support old DP and REPORT macros with printf syntax
#define FORMAT_TO_STR(Format, ...) \
- ::llvm::omptarget::debug::formatToStr(Format __VA_OPT__(, ) __VA_ARGS__)
+ ::llvm::omp::target::debug::formatToStr(Format __VA_OPT__(, ) __VA_ARGS__)
#define DP(...) ODBG() << FORMAT_TO_STR(__VA_ARGS__);
#define REPORT_INT_OLD(...) \
do { \
if (::llvm::offload::debug::isDebugEnabled()) { \
- ODBG(::llvm::omptarget::debug::ODT_Error, \
- ::llvm::omptarget::debug::ODL_Error) \
+ ODBG(::llvm::omp::target::debug::ODT_Error, \
+ ::llvm::omp::target::debug::ODL_Error) \
<< FORMAT_TO_STR(__VA_ARGS__); \
} else { \
FAILURE_MESSAGE(__VA_ARGS__); \
@@ -533,11 +533,11 @@ static inline odbg_ostream reportErrorStream() {
#endif // OMPTARGET_DEBUG
// This is used for the new style REPORT macro
-#define REPORT_INT() ::llvm::omptarget::debug::reportErrorStream()
+#define REPORT_INT() ::llvm::omp::target::debug::reportErrorStream()
// Make REPORT compatible with old and new syntax
#define REPORT(...) REPORT_INT##__VA_OPT__(_OLD)(__VA_ARGS__)
-} // namespace llvm::omptarget::debug
+} // namespace llvm::omp::target::debug
#endif // OMPTARGET_SHARED_DEBUG_H
diff --git a/offload/libomptarget/OffloadRTL.cpp b/offload/libomptarget/OffloadRTL.cpp
index 3dc37db1e1d67..3a18d76aaae15 100644
--- a/offload/libomptarget/OffloadRTL.cpp
+++ b/offload/libomptarget/OffloadRTL.cpp
@@ -19,7 +19,7 @@
#ifdef OMPT_SUPPORT
extern void llvm::omp::target::ompt::connectLibrary();
#endif
-using namespace llvm::omptarget::debug;
+using namespace llvm::omp::target::debug;
static std::mutex PluginMtx;
static uint32_t RefCount = 0;
diff --git a/offload/libomptarget/device.cpp b/offload/libomptarget/device.cpp
index 055b901372a37..e5434f68c2105 100644
--- a/offload/libomptarget/device.cpp
+++ b/offload/libomptarget/device.cpp
@@ -38,7 +38,7 @@ using namespace llvm::omp::target::ompt;
#endif
using namespace llvm::omp::target::plugin;
-using namespace llvm::omptarget::debug;
+using namespace llvm::omp::target::debug;
int HostDataToTargetTy::addEventIfNecessary(DeviceTy &Device,
AsyncInfoTy &AsyncInfo) const {
>From 35f147a08988ce800c58ba716e5f5a39c1d92406 Mon Sep 17 00:00:00 2001
From: Alex Duran <alejandro.duran at intel.com>
Date: Wed, 3 Dec 2025 06:16:18 +0100
Subject: [PATCH 4/7] [OFFLOAD][LIBOMPTARGET] More debug messages updated
---
offload/include/Shared/Debug.h | 62 +++++++++++++++
offload/libomptarget/interface.cpp | 85 +++++++++++----------
offload/libomptarget/omptarget.cpp | 117 +++++++++++++++--------------
3 files changed, 168 insertions(+), 96 deletions(-)
diff --git a/offload/include/Shared/Debug.h b/offload/include/Shared/Debug.h
index 9e657e64484c0..9fc03a0183016 100644
--- a/offload/include/Shared/Debug.h
+++ b/offload/include/Shared/Debug.h
@@ -430,6 +430,60 @@ static inline raw_ostream &operator<<(raw_ostream &Os,
#define ODBG_RESET_LEVEL() \
static_cast<llvm::offload::debug::odbg_ostream::IfLevel>(0)
+// helper templates to support lambdas with different number of arguments
+
+template <typename LambdaTy> struct lambdaHelper {
+ template <typename FuncTy, typename RetTy, typename... Args>
+ static constexpr size_t CountArgs(RetTy (FuncTy::*)(Args...)) {
+ return sizeof...(Args);
+ }
+
+ template <typename FuncTy, typename RetTy, typename... Args>
+ static constexpr size_t CountArgs(RetTy (FuncTy::*)(Args...) const) {
+ return sizeof...(Args);
+ }
+
+ static constexpr size_t NArgs = CountArgs(&LambdaTy::operator());
+
+ static void dispatch(LambdaTy func, llvm::raw_ostream &Os, uint32_t Level) {
+ if constexpr (NArgs == 1)
+ func(Os);
+ else if constexpr (NArgs == 2)
+ func(Os, Level);
+ else
+ static_assert(true, "Unsupported number of arguments in debug callback");
+ }
+};
+
+#define ODBG_OS_BASE(Stream, Component, Prefix, Type, Level, Callback) \
+ if (::llvm::offload::debug::isDebugEnabled()) { \
+ uint32_t RealLevel = (Level); \
+ if (::llvm::offload::debug::shouldPrintDebug((Component), (Type), \
+ RealLevel)) { \
+ ::llvm::offload::debug::odbg_ostream OS{ \
+ ::llvm::offload::debug::computePrefix((Prefix), (Type)), (Stream), \
+ RealLevel, /*ShouldPrefixNextString=*/true, \
+ /*ShouldEmitNewLineOnDestruction=*/true}; \
+ auto F = Callback; \
+ ::llvm::offload::debug::lambdaHelper<decltype(F)>::dispatch(F, OS, \
+ RealLevel); \
+ } \
+ }
+
+#define ODBG_OS_STREAM(Stream, Type, Level, Callback) \
+ ODBG_OS_BASE(Stream, GETNAME(TARGET_NAME), DEBUG_PREFIX, Type, Level, \
+ Callback)
+#define ODBG_OS_3(Type, Level, Callback) \
+ ODBG_OS_STREAM(llvm::offload::debug::dbgs(), Type, Level, Callback)
+#define ODBG_OS_2(Type, Callback) ODBG_OS_3(Type, 1, Callback)
+#define ODBG_OS_1(Callback) ODBG_OS_2("default", Callback)
+#define ODBG_OS_SELECT(Type, Level, Callback, NArgs, ...) ODBG_OS_##NArgs
+// Print a debug message of a certain type and verbosity level using a callback
+// to emit the message. If no type or level is provided, "default" and "1 are
+// assumed respectively.
+#define ODBG_OS(...) \
+ ODBG_OS_SELECT(__VA_ARGS__ __VA_OPT__(, ) 3, 2, 1)(__VA_ARGS__)
+
#else
inline bool isDebugEnabled() { return false; }
@@ -446,6 +500,10 @@ inline bool isDebugEnabled() { return false; }
#define ODBG_RESET_LEVEL() 0
#define ODBG(...) ODBG_NULL
+#define ODBG_OS_BASE(Stream, Component, Prefix, Type, Level, Callback)
+#define ODBG_OS_STREAM(Stream, Type, Level, Callback)
+#define ODBG_OS(...)
+
#endif
} // namespace llvm::offload::debug
@@ -476,6 +534,8 @@ constexpr const char *ODT_DumpTable = "DumpTable";
constexpr const char *ODT_MappingChanged = "MappingChanged";
constexpr const char *ODT_PluginKernel = "PluginKernel";
constexpr const char *ODT_EmptyMapping = "EmptyMapping";
+constexpr const char *ODT_Device = "Device";
+constexpr const char *ODT_Interface = "Interface";
static inline odbg_ostream reportErrorStream() {
#ifdef OMPTARGET_DEBUG
@@ -540,4 +600,6 @@ static inline odbg_ostream reportErrorStream() {
} // namespace llvm::omp::target::debug
+inline int getDebugLevel() { return 1; }
+
#endif // OMPTARGET_SHARED_DEBUG_H
diff --git a/offload/libomptarget/interface.cpp b/offload/libomptarget/interface.cpp
index fe18289765906..c17e3e39b04b9 100644
--- a/offload/libomptarget/interface.cpp
+++ b/offload/libomptarget/interface.cpp
@@ -25,6 +25,7 @@
#include "Utils/ExponentialBackoff.h"
#include "llvm/Frontend/OpenMP/OMPConstants.h"
+#include "llvm/Support/Format.h"
#include <cassert>
#include <cstdint>
@@ -35,6 +36,7 @@
#ifdef OMPT_SUPPORT
using namespace llvm::omp::target::ompt;
#endif
+using namespace llvm::omp::target::debug;
// If offload is enabled, ensure that device DeviceID has been initialized.
//
@@ -49,25 +51,25 @@ using namespace llvm::omp::target::ompt;
// This step might be skipped if offload is disabled.
bool checkDevice(int64_t &DeviceID, ident_t *Loc) {
if (OffloadPolicy::get(*PM).Kind == OffloadPolicy::DISABLED) {
- DP("Offload is disabled\n");
+ ODBG(ODT_Device) << "Offload is disabled";
return true;
}
if (DeviceID == OFFLOAD_DEVICE_DEFAULT) {
DeviceID = omp_get_default_device();
- DP("Use default device id %" PRId64 "\n", DeviceID);
+ ODBG(ODT_Device) << "Use default device id " << DeviceID;
}
// Proposed behavior for OpenMP 5.2 in OpenMP spec github issue 2669.
if (omp_get_num_devices() == 0) {
- DP("omp_get_num_devices() == 0 but offload is manadatory\n");
+ ODBG(ODT_Device) << "omp_get_num_devices() == 0 but offload is manadatory";
handleTargetOutcome(false, Loc);
return true;
}
if (DeviceID == omp_get_initial_device()) {
- DP("Device is host (%" PRId64 "), returning as if offload is disabled\n",
- DeviceID);
+ ODBG(ODT_Device) << "Device is host (" << DeviceID
+ << "), returning as if offload is disabled";
return true;
}
return false;
@@ -123,25 +125,25 @@ targetData(ident_t *Loc, int64_t DeviceId, int32_t ArgNum, void **ArgsBase,
TIMESCOPE_WITH_DETAILS_AND_IDENT("Runtime: Data Copy",
"NumArgs=" + std::to_string(ArgNum), Loc);
- DP("Entering data %s region for device %" PRId64 " with %d mappings\n",
- RegionName, DeviceId, ArgNum);
+ ODBG(ODT_Interface) << "Entering data " << RegionName << " region for device "
+ << DeviceId << " with " << ArgNum << " mappings";
if (checkDevice(DeviceId, Loc)) {
- DP("Not offloading to device %" PRId64 "\n", DeviceId);
+ ODBG(ODT_Interface) << "Not offloading to device " << DeviceId;
return;
}
if (getInfoLevel() & OMP_INFOTYPE_KERNEL_ARGS)
printKernelArguments(Loc, DeviceId, ArgNum, ArgSizes, ArgTypes, ArgNames,
RegionTypeMsg);
-#ifdef OMPTARGET_DEBUG
- for (int I = 0; I < ArgNum; ++I) {
- DP("Entry %2d: Base=" DPxMOD ", Begin=" DPxMOD ", Size=%" PRId64
- ", Type=0x%" PRIx64 ", Name=%s\n",
- I, DPxPTR(ArgsBase[I]), DPxPTR(Args[I]), ArgSizes[I], ArgTypes[I],
- (ArgNames) ? getNameFromMapping(ArgNames[I]).c_str() : "unknown");
- }
-#endif
+ ODBG_OS(ODT_Kernel, [&](llvm::raw_ostream &Os) {
+ for (int I = 0; I < ArgNum; ++I) {
+ Os << "Entry " << llvm::format_decimal(I, 2) << ": Base=" << ArgsBase[I]
+ << ", Begin=" << Args[I] << ", Size=" << ArgSizes[I]
+ << ", Type=" << llvm::format_hex(ArgTypes[I], 8) << ", Name="
+ << ((ArgNames) ? getNameFromMapping(ArgNames[I]) : "unknown") << "\n";
+ }
+ });
auto DeviceOrErr = PM->getDevice(DeviceId);
if (!DeviceOrErr)
@@ -274,7 +276,7 @@ static KernelArgsTy *upgradeKernelArgs(KernelArgsTy *KernelArgs,
KernelArgsTy &LocalKernelArgs,
int32_t NumTeams, int32_t ThreadLimit) {
if (KernelArgs->Version > OMP_KERNEL_ARG_VERSION)
- DP("Unexpected ABI version: %u\n", KernelArgs->Version);
+ ODBG(ODT_Interface) << "Unexpected ABI version: " << KernelArgs->Version;
uint32_t UpgradedVersion = KernelArgs->Version;
if (KernelArgs->Version < OMP_KERNEL_ARG_VERSION) {
@@ -326,12 +328,11 @@ static inline int targetKernel(ident_t *Loc, int64_t DeviceId, int32_t NumTeams,
assert(PM && "Runtime not initialized");
static_assert(std::is_convertible_v<TargetAsyncInfoTy &, AsyncInfoTy &>,
"Target AsyncInfoTy must be convertible to AsyncInfoTy.");
- DP("Entering target region for device %" PRId64 " with entry point " DPxMOD
- "\n",
- DeviceId, DPxPTR(HostPtr));
+ ODBG(ODT_Interface) << "Entering target region for device " << DeviceId
+ << " with entry point " << HostPtr;
if (checkDevice(DeviceId, Loc)) {
- DP("Not offloading to device %" PRId64 "\n", DeviceId);
+ ODBG(ODT_Interface) << "Not offloading to device " << DeviceId;
return OMP_TGT_FAIL;
}
@@ -354,17 +355,21 @@ static inline int targetKernel(ident_t *Loc, int64_t DeviceId, int32_t NumTeams,
printKernelArguments(Loc, DeviceId, KernelArgs->NumArgs,
KernelArgs->ArgSizes, KernelArgs->ArgTypes,
KernelArgs->ArgNames, "Entering OpenMP kernel");
-#ifdef OMPTARGET_DEBUG
- for (uint32_t I = 0; I < KernelArgs->NumArgs; ++I) {
- DP("Entry %2d: Base=" DPxMOD ", Begin=" DPxMOD ", Size=%" PRId64
- ", Type=0x%" PRIx64 ", Name=%s\n",
- I, DPxPTR(KernelArgs->ArgBasePtrs[I]), DPxPTR(KernelArgs->ArgPtrs[I]),
- KernelArgs->ArgSizes[I], KernelArgs->ArgTypes[I],
- (KernelArgs->ArgNames)
- ? getNameFromMapping(KernelArgs->ArgNames[I]).c_str()
- : "unknown");
- }
-#endif
+
+ ODBG_OS(ODT_Kernel, [&](llvm::raw_ostream &Os) {
+ for (uint32_t I = 0; I < KernelArgs->NumArgs; ++I) {
+ Os << "Entry " << llvm::format_decimal(I, 2)
+ << " Base=" << KernelArgs->ArgBasePtrs[I]
+ << ", Begin=" << KernelArgs->ArgPtrs[I]
+ << ", Size=" << KernelArgs->ArgSizes[I]
+ << ", Type=" << llvm::format_hex(KernelArgs->ArgTypes[I], 8)
+ << ", Name="
+ << (KernelArgs->ArgNames
+ ? getNameFromMapping(KernelArgs->ArgNames[I]).c_str()
+ : "unknown")
+ << "\n";
+ }
+ });
auto DeviceOrErr = PM->getDevice(DeviceId);
if (!DeviceOrErr)
@@ -463,7 +468,7 @@ EXTERN int __tgt_target_kernel_replay(ident_t *Loc, int64_t DeviceId,
assert(PM && "Runtime not initialized");
OMPT_IF_BUILT(ReturnAddressSetterRAII RA(__builtin_return_address(0)));
if (checkDevice(DeviceId, Loc)) {
- DP("Not offloading to device %" PRId64 "\n", DeviceId);
+ ODBG(ODT_Interface) << "Not offloading to device " << DeviceId;
return OMP_TGT_FAIL;
}
auto DeviceOrErr = PM->getDevice(DeviceId);
@@ -491,8 +496,8 @@ EXTERN int __tgt_target_kernel_replay(ident_t *Loc, int64_t DeviceId,
EXTERN int64_t __tgt_mapper_num_components(void *RtMapperHandle) {
auto *MapperComponentsPtr = (struct MapperComponentsTy *)RtMapperHandle;
int64_t Size = MapperComponentsPtr->Components.size();
- DP("__tgt_mapper_num_components(Handle=" DPxMOD ") returns %" PRId64 "\n",
- DPxPTR(RtMapperHandle), Size);
+ ODBG(ODT_Interface) << "__tgt_mapper_num_components(Handle=" << RtMapperHandle
+ << ") returns " << Size;
return Size;
}
@@ -500,11 +505,11 @@ EXTERN int64_t __tgt_mapper_num_components(void *RtMapperHandle) {
EXTERN void __tgt_push_mapper_component(void *RtMapperHandle, void *Base,
void *Begin, int64_t Size, int64_t Type,
void *Name) {
- DP("__tgt_push_mapper_component(Handle=" DPxMOD
- ") adds an entry (Base=" DPxMOD ", Begin=" DPxMOD ", Size=%" PRId64
- ", Type=0x%" PRIx64 ", Name=%s).\n",
- DPxPTR(RtMapperHandle), DPxPTR(Base), DPxPTR(Begin), Size, Type,
- (Name) ? getNameFromMapping(Name).c_str() : "unknown");
+ ODBG(ODT_Interface) << "__tgt_push_mapper_component(Handle=" << RtMapperHandle
+ << ") adds an entry (Base=" << Base << ", Begin=" << Begin
+ << ", Size=" << Size
+ << ", Type=" << llvm::format_hex(Type, 8) << ", Name="
+ << ((Name) ? getNameFromMapping(Name) : "unknown") << ")";
auto *MapperComponentsPtr = (struct MapperComponentsTy *)RtMapperHandle;
MapperComponentsPtr->Components.push_back(
MapComponentInfoTy(Base, Begin, Size, Type, Name));
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index 69725e77bae00..e0ff7834afce3 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -41,6 +41,7 @@ using llvm::SmallVector;
#ifdef OMPT_SUPPORT
using namespace llvm::omp::target::ompt;
#endif
+using namespace llvm::omp::target::debug;
int AsyncInfoTy::synchronize() {
int Result = OFFLOAD_SUCCESS;
@@ -200,10 +201,11 @@ static int32_t getParentIndex(int64_t Type) {
void *targetAllocExplicit(size_t Size, int DeviceNum, int Kind,
const char *Name) {
- DP("Call to %s for device %d requesting %zu bytes\n", Name, DeviceNum, Size);
+ ODBG(ODT_Interface) << "Call to " << Name << " for device " << DeviceNum
+ << " requesting " << Size << " bytes";
if (Size <= 0) {
- DP("Call to %s with non-positive length\n", Name);
+ ODBG(ODT_Interface) << "Call to " << Name << " with non-positive length";
return NULL;
}
@@ -211,7 +213,7 @@ void *targetAllocExplicit(size_t Size, int DeviceNum, int Kind,
if (DeviceNum == omp_get_initial_device()) {
Rc = malloc(Size);
- DP("%s returns host ptr " DPxMOD "\n", Name, DPxPTR(Rc));
+ ODBG(ODT_Interface) << Name << " returns host ptr " << Rc;
return Rc;
}
@@ -220,23 +222,23 @@ void *targetAllocExplicit(size_t Size, int DeviceNum, int Kind,
FATAL_MESSAGE(DeviceNum, "%s", toString(DeviceOrErr.takeError()).c_str());
Rc = DeviceOrErr->allocData(Size, nullptr, Kind);
- DP("%s returns device ptr " DPxMOD "\n", Name, DPxPTR(Rc));
+ ODBG(ODT_Interface) << Name << " returns device ptr " << Rc;
return Rc;
}
void targetFreeExplicit(void *DevicePtr, int DeviceNum, int Kind,
const char *Name) {
- DP("Call to %s for device %d and address " DPxMOD "\n", Name, DeviceNum,
- DPxPTR(DevicePtr));
+ ODBG(ODT_Interface) << "Call to " << Name << " for device " << DeviceNum
+ << " and address " << DevicePtr;
if (!DevicePtr) {
- DP("Call to %s with NULL ptr\n", Name);
+ ODBG(ODT_Interface) << "Call to " << Name << " with NULL ptr";
return;
}
if (DeviceNum == omp_get_initial_device()) {
free(DevicePtr);
- DP("%s deallocated host ptr\n", Name);
+ ODBG(ODT_Interface) << Name << " deallocated host ptr";
return;
}
@@ -249,15 +251,16 @@ void targetFreeExplicit(void *DevicePtr, int DeviceNum, int Kind,
"Failed to deallocate device ptr. Set "
"OFFLOAD_TRACK_ALLOCATION_TRACES=1 to track allocations.");
- DP("omp_target_free deallocated device ptr\n");
+ ODBG(ODT_Interface) << "omp_target_free deallocated device ptr";
}
void *targetLockExplicit(void *HostPtr, size_t Size, int DeviceNum,
const char *Name) {
- DP("Call to %s for device %d locking %zu bytes\n", Name, DeviceNum, Size);
+ ODBG(ODT_Interface) << "Call to " << Name << " for device " << DeviceNum
+ << " locking " << Size << " bytes";
if (Size <= 0) {
- DP("Call to %s with non-positive length\n", Name);
+ ODBG(ODT_Interface) << "Call to " << Name << " with non-positive length";
return NULL;
}
@@ -270,22 +273,23 @@ void *targetLockExplicit(void *HostPtr, size_t Size, int DeviceNum,
int32_t Err = 0;
Err = DeviceOrErr->RTL->data_lock(DeviceNum, HostPtr, Size, &RC);
if (Err) {
- DP("Could not lock ptr %p\n", HostPtr);
+ ODBG(ODT_Interface) << "Could not lock ptr " << HostPtr;
return nullptr;
}
- DP("%s returns device ptr " DPxMOD "\n", Name, DPxPTR(RC));
+ ODBG(ODT_Interface) << Name << " returns device ptr " << RC;
return RC;
}
void targetUnlockExplicit(void *HostPtr, int DeviceNum, const char *Name) {
- DP("Call to %s for device %d unlocking\n", Name, DeviceNum);
+ ODBG(ODT_Interface) << "Call to " << Name << " for device " << DeviceNum
+ << " unlocking";
auto DeviceOrErr = PM->getDevice(DeviceNum);
if (!DeviceOrErr)
FATAL_MESSAGE(DeviceNum, "%s", toString(DeviceOrErr.takeError()).c_str());
DeviceOrErr->RTL->data_unlock(DeviceNum, HostPtr);
- DP("%s returns\n", Name);
+ ODBG(ODT_Interface) << Name << " returns";
}
/// Call the user-defined mapper function followed by the appropriate
@@ -295,7 +299,7 @@ int targetDataMapper(ident_t *Loc, DeviceTy &Device, void *ArgBase, void *Arg,
void *ArgMapper, AsyncInfoTy &AsyncInfo,
TargetDataFuncPtrTy TargetDataFunction,
AttachInfoTy *AttachInfo = nullptr) {
- DP("Calling the mapper function " DPxMOD "\n", DPxPTR(ArgMapper));
+ ODBG(ODT_Interface) << "Calling the mapper function " << ArgMapper;
// The mapper function fills up Components.
MapperComponentsTy MapperComponents;
@@ -368,12 +372,11 @@ static void *calculateTargetPointeeBase(void *HstPteeBase, void *HstPteeBegin,
void *TgtPteeBase = reinterpret_cast<void *>(
reinterpret_cast<uint64_t>(TgtPteeBegin) - Delta);
- DP("HstPteeBase: " DPxMOD ", HstPteeBegin: " DPxMOD
- ", Delta (HstPteeBegin - HstPteeBase): %" PRIu64 ".\n",
- DPxPTR(HstPteeBase), DPxPTR(HstPteeBegin), Delta);
- DP("TgtPteeBase (TgtPteeBegin - Delta): " DPxMOD ", TgtPteeBegin : " DPxMOD
- "\n",
- DPxPTR(TgtPteeBase), DPxPTR(TgtPteeBegin));
+ ODBG(ODT_Mapping) << "HstPteeBase: " << HstPteeBase
+ << ", HstPteeBegin: " << HstPteeBegin
+ << ", Delta (HstPteeBegin - HstPteeBase): " << Delta << "\n"
+ << "TgtPteeBase (TgtPteeBegin - Delta): " << TgtPteeBase
+ << ", TgtPteeBegin: " << TgtPteeBegin;
return TgtPteeBase;
}
@@ -453,8 +456,8 @@ static int performPointerAttachment(DeviceTy &Device, AsyncInfoTy &AsyncInfo,
// Add shadow pointer tracking
if (!PtrTPR.getEntry()->addShadowPointer(
ShadowPtrInfoTy{HstPtrAddr, TgtPtrAddr, TgtPteeBase, HstPtrSize})) {
- DP("Pointer " DPxMOD " is already attached to " DPxMOD "\n",
- DPxPTR(TgtPtrAddr), DPxPTR(TgtPteeBase));
+ ODBG(ODT_Mapping) << "Pointer " << TgtPtrAddr << " is already attached to "
+ << TgtPteeBase;
return OFFLOAD_SUCCESS;
}
@@ -464,7 +467,7 @@ static int performPointerAttachment(DeviceTy &Device, AsyncInfoTy &AsyncInfo,
// Lambda to handle submitData result and perform final steps.
auto HandleSubmitResult = [&](int SubmitResult) -> int {
if (SubmitResult != OFFLOAD_SUCCESS) {
- REPORT("Failed to update pointer on device.\n");
+ REPORT() << "Failed to update pointer on device.";
return OFFLOAD_FAIL;
}
@@ -532,8 +535,8 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
targetDataBegin, AttachInfo);
if (Rc != OFFLOAD_SUCCESS) {
- REPORT("Call to targetDataBegin via targetDataMapper for custom mapper"
- " failed.\n");
+ REPORT() << "Call to targetDataBegin via targetDataMapper for custom "
+ "mapper failed";
return OFFLOAD_FAIL;
}
@@ -575,9 +578,8 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
int64_t Alignment = getPartialStructRequiredAlignment(HstPtrBase);
TgtPadding = (int64_t)HstPtrBegin % Alignment;
if (TgtPadding) {
- DP("Using a padding of %" PRId64 " bytes for begin address " DPxMOD
- "\n",
- TgtPadding, DPxPTR(HstPtrBegin));
+ ODBG(ODT_Mapping) << "Using a padding of " << TgtPadding
+ << " bytes for begin address " << HstPtrBegin;
}
}
@@ -602,7 +604,7 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
MappingInfoTy::HDTTMapAccessorTy HDTTMap =
Device.getMappingInfo().HostDataToTargetMap.getExclusiveAccessor();
if (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ) {
- DP("Has a pointer entry: \n");
+ ODBG(ODT_Mapping) << "Has a pointer entry";
// Base is address of pointer.
//
// Usually, the pointer is already allocated by this time. For example:
@@ -625,9 +627,10 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
PointerTgtPtrBegin = PointerTpr.TargetPointer;
IsHostPtr = PointerTpr.Flags.IsHostPointer;
if (!PointerTgtPtrBegin) {
- REPORT("Call to getTargetPointer returned null pointer (%s).\n",
- HasPresentModifier ? "'present' map type modifier"
- : "device failure or illegal mapping");
+ REPORT() << "Call to getTargetPointer returned null pointer ("
+ << (HasPresentModifier ? "'present' map type modifier"
+ : "device failure or illegal mapping")
+ << ")";
return OFFLOAD_FAIL;
}
@@ -660,9 +663,10 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
// If data_size==0, then the argument could be a zero-length pointer to
// NULL, so getOrAlloc() returning NULL is not an error.
if (!TgtPtrBegin && (DataSize || HasPresentModifier)) {
- REPORT("Call to getTargetPointer returned null pointer (%s).\n",
- HasPresentModifier ? "'present' map type modifier"
- : "device failure or illegal mapping");
+ REPORT() << "Call to getTargetPointer returned null pointer ("
+ << (HasPresentModifier ? "'present' map type modifier"
+ : "device failure or illegal mapping")
+ << ").";
return OFFLOAD_FAIL;
}
@@ -868,7 +872,7 @@ int processAttachEntries(DeviceTy &Device, AttachInfoTy &AttachInfo,
DP("Inserting a data fence before the first pointer attachment.\n");
Ret = Device.dataFence(AsyncInfo);
if (Ret != OFFLOAD_SUCCESS) {
- REPORT("Failed to insert data fence.\n");
+ REPORT() << "Failed to insert data fence.";
return OFFLOAD_FAIL;
}
}
@@ -1040,8 +1044,8 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
targetDataEnd);
if (Ret != OFFLOAD_SUCCESS) {
- REPORT("Call to targetDataEnd via targetDataMapper for custom mapper"
- " failed.\n");
+ REPORT() << "Call to targetDataEnd via targetDataMapper for custom "
+ "mapper failed.";
return OFFLOAD_FAIL;
}
@@ -1123,7 +1127,7 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
Ret = Device.retrieveData(HstPtrBegin, TgtPtrBegin, DataSize, AsyncInfo,
TPR.getEntry());
if (Ret != OFFLOAD_SUCCESS) {
- REPORT("Copying data from device failed.\n");
+ REPORT() << "Copying data from device failed.";
return OFFLOAD_FAIL;
}
@@ -1185,7 +1189,7 @@ static int targetDataContiguous(ident_t *Loc, DeviceTy &Device, void *ArgsBase,
int Ret = Device.submitData(TgtPtrBegin, HstPtrBegin, ArgSize, AsyncInfo,
TPR.getEntry());
if (Ret != OFFLOAD_SUCCESS) {
- REPORT("Copying data to device failed.\n");
+ REPORT() << "Copying data to device failed.";
return OFFLOAD_FAIL;
}
if (TPR.getEntry()) {
@@ -1208,7 +1212,7 @@ static int targetDataContiguous(ident_t *Loc, DeviceTy &Device, void *ArgsBase,
ShadowPtr.TgtPtrContent.data(),
ShadowPtr.PtrSize, AsyncInfo);
if (Ret != OFFLOAD_SUCCESS) {
- REPORT("Copying data to device failed.\n");
+ REPORT() << "Copying data to device failed.";
return OFFLOAD_FAIL;
}
return OFFLOAD_SUCCESS;
@@ -1226,7 +1230,7 @@ static int targetDataContiguous(ident_t *Loc, DeviceTy &Device, void *ArgsBase,
int Ret = Device.retrieveData(HstPtrBegin, TgtPtrBegin, ArgSize, AsyncInfo,
TPR.getEntry());
if (Ret != OFFLOAD_SUCCESS) {
- REPORT("Copying data from device failed.\n");
+ REPORT() << "Copying data from device failed.";
return OFFLOAD_FAIL;
}
@@ -1334,8 +1338,8 @@ int targetDataUpdate(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
targetDataUpdate);
if (Ret != OFFLOAD_SUCCESS) {
- REPORT("Call to targetDataUpdate via targetDataMapper for custom mapper"
- " failed.\n");
+ REPORT() << "Call to targetDataUpdate via targetDataMapper for custom "
+ "mapper failed.";
return OFFLOAD_FAIL;
}
@@ -1814,7 +1818,7 @@ static int processDataBefore(ident_t *Loc, int64_t DeviceId, void *HostPtr,
ArgTypes, ArgNames, ArgMappers, AsyncInfo,
&AttachInfo, false /*FromMapper=*/);
if (Ret != OFFLOAD_SUCCESS) {
- REPORT("Call to targetDataBegin failed, abort target.\n");
+ REPORT() << "Call to targetDataBegin failed, abort target.";
return OFFLOAD_FAIL;
}
@@ -1822,7 +1826,7 @@ static int processDataBefore(ident_t *Loc, int64_t DeviceId, void *HostPtr,
if (!AttachInfo.AttachEntries.empty()) {
Ret = processAttachEntries(*DeviceOrErr, AttachInfo, AsyncInfo);
if (Ret != OFFLOAD_SUCCESS) {
- REPORT("Failed to process ATTACH entries.\n");
+ REPORT() << "Failed to process ATTACH entries.";
return OFFLOAD_FAIL;
}
}
@@ -1873,7 +1877,7 @@ static int processDataBefore(ident_t *Loc, int64_t DeviceId, void *HostPtr,
DeviceOrErr->submitData(TgtPtrBegin, &PointerTgtPtrBegin,
sizeof(void *), AsyncInfo, TPR.getEntry());
if (Ret != OFFLOAD_SUCCESS) {
- REPORT("Copying data to device failed.\n");
+ REPORT() << "Copying data to device failed.";
return OFFLOAD_FAIL;
}
}
@@ -1936,9 +1940,10 @@ static int processDataBefore(ident_t *Loc, int64_t DeviceId, void *HostPtr,
/*TgtArgsIndex=*/TgtArgs.size(), HstPtrName, AllocImmediately,
HstPteeBase, HstPteeBegin, /*IsCorrespondingPointerInit=*/IsAttach);
if (Ret != OFFLOAD_SUCCESS) {
- REPORT("Failed to process %s%sprivate argument " DPxMOD "\n",
- IsAttach ? "corresponding-pointer-initialization " : "",
- (IsFirstPrivate ? "first-" : ""), DPxPTR(HstPtrBegin));
+ REPORT() << "Failed to process "
+ << (IsAttach ? "corresponding-pointer-initialization " : "")
+ << (IsFirstPrivate ? "first-" : "") << "private argument "
+ << HstPtrBegin << ".";
return OFFLOAD_FAIL;
}
} else {
@@ -1991,7 +1996,7 @@ static int processDataAfter(ident_t *Loc, int64_t DeviceId, void *HostPtr,
int Ret = targetDataEnd(Loc, *DeviceOrErr, ArgNum, ArgBases, Args, ArgSizes,
ArgTypes, ArgNames, ArgMappers, AsyncInfo);
if (Ret != OFFLOAD_SUCCESS) {
- REPORT("Call to targetDataEnd failed, abort target.\n");
+ REPORT() << "Call to targetDataEnd failed, abort target.";
return OFFLOAD_FAIL;
}
@@ -2003,7 +2008,7 @@ static int processDataAfter(ident_t *Loc, int64_t DeviceId, void *HostPtr,
std::move(PrivateArgumentManager)]() mutable -> int {
int Ret = PrivateArgumentManager.free();
if (Ret != OFFLOAD_SUCCESS) {
- REPORT("Failed to deallocate target memory for private args\n");
+ REPORT() << "Failed to deallocate target memory for private args";
return OFFLOAD_FAIL;
}
return Ret;
@@ -2066,7 +2071,7 @@ int target(ident_t *Loc, DeviceTy &Device, void *HostPtr,
KernelArgs.ArgNames, KernelArgs.ArgMappers, TgtArgs,
TgtOffsets, PrivateArgumentManager, AsyncInfo);
if (Ret != OFFLOAD_SUCCESS) {
- REPORT("Failed to process data before launching the kernel.\n");
+ REPORT() << "Failed to process data before launching the kernel.";
return OFFLOAD_FAIL;
}
@@ -2118,7 +2123,7 @@ int target(ident_t *Loc, DeviceTy &Device, void *HostPtr,
KernelArgs.ArgNames, KernelArgs.ArgMappers,
PrivateArgumentManager, AsyncInfo);
if (Ret != OFFLOAD_SUCCESS) {
- REPORT("Failed to process data after launching the kernel.\n");
+ REPORT() << "Failed to process data after launching the kernel.";
return OFFLOAD_FAIL;
}
}
>From 77d8bf877237f97c652ef22ce91ba6d75c6e58c2 Mon Sep 17 00:00:00 2001
From: Alex Duran <alejandro.duran at intel.com>
Date: Wed, 3 Dec 2025 06:52:00 +0100
Subject: [PATCH 5/7] more messages in omptarget.cpp
---
offload/include/Shared/Debug.h | 1 +
offload/libomptarget/omptarget.cpp | 334 +++++++++++++++--------------
2 files changed, 178 insertions(+), 157 deletions(-)
diff --git a/offload/include/Shared/Debug.h b/offload/include/Shared/Debug.h
index 9fc03a0183016..d60ba26257bbe 100644
--- a/offload/include/Shared/Debug.h
+++ b/offload/include/Shared/Debug.h
@@ -536,6 +536,7 @@ constexpr const char *ODT_PluginKernel = "PluginKernel";
constexpr const char *ODT_EmptyMapping = "EmptyMapping";
constexpr const char *ODT_Device = "Device";
constexpr const char *ODT_Interface = "Interface";
+constexpr const char *ODT_Alloc = "Alloc";
static inline odbg_ostream reportErrorStream() {
#ifdef OMPTARGET_DEBUG
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index e0ff7834afce3..021caff159919 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -461,8 +461,8 @@ static int performPointerAttachment(DeviceTy &Device, AsyncInfoTy &AsyncInfo,
return OFFLOAD_SUCCESS;
}
- DP("Update pointer (" DPxMOD ") -> [" DPxMOD "]\n", DPxPTR(TgtPtrAddr),
- DPxPTR(TgtPteeBase));
+ ODBG(ODT_Mapping) << "Update pointer (" << TgtPtrAddr << ") -> ["
+ << TgtPteeBase << "]\n";
// Lambda to handle submitData result and perform final steps.
auto HandleSubmitResult = [&](int SubmitResult) -> int {
@@ -494,11 +494,11 @@ static int performPointerAttachment(DeviceTy &Device, AsyncInfoTy &AsyncInfo,
std::memcpy(SrcBuffer + VoidPtrSize, HstDescriptorFieldsAddr,
HstDescriptorFieldsSize);
- DP("Updating %" PRId64 " bytes of descriptor (" DPxMOD
- ") (pointer + %" PRId64 " additional bytes from host descriptor " DPxMOD
- ")\n",
- HstPtrSize, DPxPTR(TgtPtrAddr), HstDescriptorFieldsSize,
- DPxPTR(HstDescriptorFieldsAddr));
+ ODBG(ODT_Mapping) << "Updating " << HstPtrSize << " bytes of descriptor ("
+ << TgtPtrAddr << ") (pointer + "
+ << HstDescriptorFieldsSize
+ << " additional bytes from host descriptor "
+ << HstDescriptorFieldsAddr << ")";
}
// Submit the populated source buffer to device.
@@ -527,7 +527,8 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
// Instead of executing the regular path of targetDataBegin, call the
// targetDataMapper variant which will call targetDataBegin again
// with new arguments.
- DP("Calling targetDataMapper for the %dth argument\n", I);
+ ODBG(ODT_Mapping) << "Calling targetDataMapper for the " << I
+ << "th argument";
map_var_info_t ArgName = (!ArgNames) ? nullptr : ArgNames[I];
int Rc = targetDataMapper(Loc, Device, ArgsBase[I], Args[I], ArgSizes[I],
@@ -564,7 +565,8 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
/*PointerSize=*/DataSize, /*MapType=*/ArgTypes[I],
/*PointeeName=*/HstPtrName);
- DP("Deferring ATTACH map-type processing for argument %d\n", I);
+ ODBG(ODT_Mapping) << "Deferring ATTACH map-type processing for argument "
+ << I;
continue;
}
@@ -638,10 +640,11 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
if (PointerTpr.Flags.IsNewEntry && !IsHostPtr)
AttachInfo->NewAllocations[HstPtrBase] = sizeof(void *);
- DP("There are %zu bytes allocated at target address " DPxMOD " - is%s new"
- "\n",
- sizeof(void *), DPxPTR(PointerTgtPtrBegin),
- (PointerTpr.Flags.IsNewEntry ? "" : " not"));
+ ODBG(ODT_Mapping) << "There are " << sizeof(void *)
+ << " bytes allocated at target address "
+ << PointerTgtPtrBegin << " - is"
+ << (PointerTpr.Flags.IsNewEntry ? "" : " not")
+ << " new";
PointerHstPtrBegin = HstPtrBase;
// modify current entry.
HstPtrBase = *reinterpret_cast<void **>(HstPtrBase);
@@ -674,14 +677,15 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
if (TPR.Flags.IsNewEntry && !IsHostPtr && TgtPtrBegin)
AttachInfo->NewAllocations[HstPtrBegin] = DataSize;
- DP("There are %" PRId64 " bytes allocated at target address " DPxMOD
- " - is%s new\n",
- DataSize, DPxPTR(TgtPtrBegin), (TPR.Flags.IsNewEntry ? "" : " not"));
+ ODBG(ODT_Mapping) << "There are " << DataSize
+ << " bytes allocated at target address " << TgtPtrBegin
+ << " - is" << (TPR.Flags.IsNewEntry ? "" : " not")
+ << " new";
if (ArgTypes[I] & OMP_TGT_MAPTYPE_RETURN_PARAM) {
uintptr_t Delta = (uintptr_t)HstPtrBegin - (uintptr_t)HstPtrBase;
void *TgtPtrBase = (void *)((uintptr_t)TgtPtrBegin - Delta);
- DP("Returning device pointer " DPxMOD "\n", DPxPTR(TgtPtrBase));
+ ODBG(ODT_Mapping) << "Returning device pointer " << TgtPtrBase;
ArgsBase[I] = TgtPtrBase;
}
@@ -759,19 +763,20 @@ int processAttachEntries(DeviceTy &Device, AttachInfoTy &AttachInfo,
AsyncInfoTy &AsyncInfo) {
// Report all tracked allocations from both main loop and ATTACH processing
if (!AttachInfo.NewAllocations.empty()) {
- DP("Tracked %u total new allocations:\n",
- (unsigned)AttachInfo.NewAllocations.size());
- for ([[maybe_unused]] const auto &Alloc : AttachInfo.NewAllocations) {
- DP(" Host ptr: " DPxMOD ", Size: %" PRId64 " bytes\n",
- DPxPTR(Alloc.first), Alloc.second);
- }
+ ODBG_OS(ODT_Mapping, [&](llvm::raw_ostream &OS) {
+ OS << "Tracked " << AttachInfo.NewAllocations.size()
+ << " total new allocations:";
+ for (const auto &Alloc : AttachInfo.NewAllocations) {
+ OS << " Host ptr: " << Alloc.first << ", Size: " << Alloc.second
+ << " bytes";
+ }
+ });
}
if (AttachInfo.AttachEntries.empty())
return OFFLOAD_SUCCESS;
- DP("Processing %zu deferred ATTACH map entries\n",
- AttachInfo.AttachEntries.size());
+ ODBG(ODT_Mapping) << "Processing " << AttachInfo.AttachEntries.size();
int Ret = OFFLOAD_SUCCESS;
bool IsFirstPointerAttachment = true;
@@ -787,9 +792,11 @@ int processAttachEntries(DeviceTy &Device, AttachInfoTy &AttachInfo,
int64_t PtrSize = AttachEntry.PointerSize;
int64_t MapType = AttachEntry.MapType;
- DP("Processing ATTACH entry %zu: HstPtr=" DPxMOD ", HstPteeBegin=" DPxMOD
- ", Size=%" PRId64 ", Type=0x%" PRIx64 "\n",
- EntryIdx, DPxPTR(HstPtr), DPxPTR(HstPteeBegin), PtrSize, MapType);
+ ODBG(ODT_Mapping) << "Processing ATTACH entry " << EntryIdx
+ << ": HstPtr=" << HstPtr
+ << ", HstPteeBegin=" << HstPteeBegin
+ << ", PtrSize=" << PtrSize << ", MapType=0x"
+ << llvm::utohexstr(MapType);
const bool IsAttachAlways = MapType & OMP_TGT_MAPTYPE_ALWAYS;
@@ -803,8 +810,9 @@ int processAttachEntries(DeviceTy &Device, AttachInfoTy &AttachInfo,
Ptr < reinterpret_cast<void *>(
reinterpret_cast<char *>(AllocPtr) + AllocSize);
});
- DP("Attach %s " DPxMOD " was newly allocated: %s\n", PtrName, DPxPTR(Ptr),
- IsNewlyAllocated ? "yes" : "no");
+ ODBG(ODT_Mapping) << "Attach " << PtrName << " " << Ptr
+ << " was newly allocated: "
+ << (IsNewlyAllocated ? "yes" : "no");
return IsNewlyAllocated;
};
@@ -812,9 +820,9 @@ int processAttachEntries(DeviceTy &Device, AttachInfoTy &AttachInfo,
// allocated, or the ALWAYS flag is set.
if (!IsAttachAlways && !WasNewlyAllocated(HstPteeBegin, "pointee") &&
!WasNewlyAllocated(HstPtr, "pointer")) {
- DP("Skipping ATTACH entry %zu: neither pointer nor pointee was newly "
- "allocated and no ALWAYS flag\n",
- EntryIdx);
+ ODBG(ODT_Mapping) << "Skipping ATTACH entry " << EntryIdx
+ << ": neither pointer nor pointee was newly "
+ "allocated and no ALWAYS flag";
continue;
}
@@ -828,19 +836,19 @@ int processAttachEntries(DeviceTy &Device, AttachInfoTy &AttachInfo,
Ptr, Size, /*UpdateRefCount=*/false,
/*UseHoldRefCount=*/false, /*MustContain=*/true);
- DP("Attach %s lookup - IsPresent=%s, IsHostPtr=%s\n", PtrType,
- TPR.isPresent() ? "yes" : "no",
- TPR.Flags.IsHostPointer ? "yes" : "no");
+ ODBG(ODT_Mapping) << "Attach " << PtrType << " lookup - IsPresent="
+ << (TPR.isPresent() ? "yes" : "no") << ", IsHostPtr="
+ << (TPR.Flags.IsHostPointer ? "yes" : "no");
if (!TPR.isPresent()) {
- DP("Skipping ATTACH entry %zu: %s not present on device\n", EntryIdx,
- PtrType);
+ ODBG(ODT_Mapping) << "Skipping ATTACH entry " << EntryIdx << ": "
+ << PtrType << " not present on device";
return std::nullopt;
}
if (TPR.Flags.IsHostPointer) {
- DP("Skipping ATTACH entry %zu: device version of the %s is a host "
- "pointer.\n",
- EntryIdx, PtrType);
+ ODBG(ODT_Mapping) << "Skipping ATTACH entry " << EntryIdx
+ << ": device version of the " << PtrType
+ << " is a host pointer.";
return std::nullopt;
}
@@ -869,7 +877,8 @@ int processAttachEntries(DeviceTy &Device, AttachInfoTy &AttachInfo,
// Insert a data-fence before the first pointer-attachment.
if (IsFirstPointerAttachment) {
IsFirstPointerAttachment = false;
- DP("Inserting a data fence before the first pointer attachment.\n");
+ ODBG(ODT_Mapping)
+ << "Inserting a data fence before the first pointer attachment.";
Ret = Device.dataFence(AsyncInfo);
if (Ret != OFFLOAD_SUCCESS) {
REPORT() << "Failed to insert data fence.";
@@ -885,7 +894,8 @@ int processAttachEntries(DeviceTy &Device, AttachInfoTy &AttachInfo,
if (Ret != OFFLOAD_SUCCESS)
return OFFLOAD_FAIL;
- DP("ATTACH entry %zu processed successfully\n", EntryIdx);
+ ODBG(ODT_Mapping) << "ATTACH entry " << EntryIdx
+ << " processed successfully";
}
return OFFLOAD_SUCCESS;
@@ -970,16 +980,16 @@ postProcessingTargetDataEnd(DeviceTy *Device,
Entry->foreachShadowPointerInfo([&](const ShadowPtrInfoTy &ShadowPtr) {
constexpr int64_t VoidPtrSize = sizeof(void *);
if (ShadowPtr.PtrSize > VoidPtrSize) {
- DP("Restoring host descriptor " DPxMOD
- " to its original content (%" PRId64
- " bytes), containing pointee address " DPxMOD "\n",
- DPxPTR(ShadowPtr.HstPtrAddr), ShadowPtr.PtrSize,
- DPxPTR(ShadowPtr.HstPtrContent.data()));
+ ODBG(ODT_Mapping)
+ << "Restoring host descriptor " << (void *)ShadowPtr.HstPtrAddr
+ << " to its original content (" << ShadowPtr.PtrSize
+ << " bytes), containing pointee address "
+ << (void *)ShadowPtr.HstPtrContent.data();
} else {
- DP("Restoring host pointer " DPxMOD " to its original value " DPxMOD
- "\n",
- DPxPTR(ShadowPtr.HstPtrAddr),
- DPxPTR(ShadowPtr.HstPtrContent.data()));
+ ODBG(ODT_Mapping)
+ << "Restoring host pointer " << (void *)ShadowPtr.HstPtrAddr
+ << " to its original value "
+ << (void *)ShadowPtr.HstPtrContent.data();
}
std::memcpy(ShadowPtr.HstPtrAddr, ShadowPtr.HstPtrContent.data(),
ShadowPtr.PtrSize);
@@ -999,7 +1009,7 @@ postProcessingTargetDataEnd(DeviceTy *Device,
HDTTMap.destroy();
Ret |= Device->getMappingInfo().deallocTgtPtrAndEntry(Entry, DataSize);
if (Ret != OFFLOAD_SUCCESS) {
- REPORT("Deallocating data from device failed.\n");
+ REPORT() << "Deallocating data from device failed.";
break;
}
}
@@ -1028,7 +1038,7 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
// directives. They may be encountered here while handling the "end" part of
// "#pragma omp target".
if (ArgTypes[I] & OMP_TGT_MAPTYPE_ATTACH) {
- DP("Ignoring ATTACH entry %d in targetDataEnd\n", I);
+ ODBG(ODT_Mapping) << "Ignoring ATTACH entry " << I << " in targetDataEnd";
continue;
}
@@ -1036,7 +1046,8 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
// Instead of executing the regular path of targetDataEnd, call the
// targetDataMapper variant which will call targetDataEnd again
// with new arguments.
- DP("Calling targetDataMapper for the %dth argument\n", I);
+ ODBG(ODT_Mapping) << "Calling targetDataMapper for the " << I
+ << "th argument";
map_var_info_t ArgName = (!ArgNames) ? nullptr : ArgNames[I];
Ret = targetDataMapper(Loc, Device, ArgBases[I], Args[I], ArgSizes[I],
@@ -1070,8 +1081,10 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
void *TgtPtrBegin = TPR.TargetPointer;
if (!TPR.isPresent() && !TPR.isHostPointer() &&
(DataSize || HasPresentModifier)) {
- DP("Mapping does not exist (%s)\n",
- (HasPresentModifier ? "'present' map type modifier" : "ignored"));
+ ODBG(ODT_Mapping) << "Mapping does not exist ("
+ << (HasPresentModifier ? "'present' map type modifier"
+ : "ignored")
+ << ")";
if (HasPresentModifier) {
// OpenMP 5.1, sec. 2.21.7.1 "map Clause", p. 350 L10-13:
// "If a map clause appears on a target, target data, target enter data
@@ -1094,9 +1107,10 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
return OFFLOAD_FAIL;
}
} else {
- DP("There are %" PRId64 " bytes allocated at target address " DPxMOD
- " - is%s last\n",
- DataSize, DPxPTR(TgtPtrBegin), (TPR.Flags.IsLast ? "" : " not"));
+ ODBG(ODT_Mapping) << "There are " << DataSize
+ << " bytes allocated at target address " << TgtPtrBegin
+ << " - is" << (TPR.Flags.IsLast ? "" : " not")
+ << " last";
}
// OpenMP 5.1, sec. 2.21.7.1 "map Clause", p. 351 L14-16:
@@ -1112,14 +1126,15 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
const bool HasFrom = ArgTypes[I] & OMP_TGT_MAPTYPE_FROM;
if (HasFrom && (HasAlways || TPR.Flags.IsLast) &&
!TPR.Flags.IsHostPointer && DataSize != 0) {
- DP("Moving %" PRId64 " bytes (tgt:" DPxMOD ") -> (hst:" DPxMOD ")\n",
- DataSize, DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin));
+ ODBG(ODT_Mapping) << "Moving " << DataSize
+ << " bytes (tgt:" << TgtPtrBegin
+ << ") -> (hst:" << HstPtrBegin << ")";
TIMESCOPE_WITH_DETAILS_AND_IDENT(
"DevToHost", "Size=" + std::to_string(DataSize) + "B", Loc);
// Wait for any previous transfer if an event is present.
if (void *Event = TPR.getEntry()->getEvent()) {
if (Device.waitEvent(Event, AsyncInfo) != OFFLOAD_SUCCESS) {
- REPORT("Failed to wait for event " DPxMOD ".\n", DPxPTR(Event));
+ REPORT() << "Failed to wait for event " << Event << ".";
return OFFLOAD_FAIL;
}
}
@@ -1167,7 +1182,8 @@ static int targetDataContiguous(ident_t *Loc, DeviceTy &Device, void *ArgsBase,
/*UseHoldRefCount=*/false, /*MustContain=*/true);
void *TgtPtrBegin = TPR.TargetPointer;
if (!TPR.isPresent()) {
- DP("hst data:" DPxMOD " not found, becomes a noop\n", DPxPTR(HstPtrBegin));
+ ODBG(ODT_Mapping) << "hst data:" << HstPtrBegin
+ << " not found, becomes a noop";
if (ArgType & OMP_TGT_MAPTYPE_PRESENT) {
MESSAGE("device mapping required by 'present' motion modifier does not "
"exist for host address " DPxMOD " (%" PRId64 " bytes)",
@@ -1178,14 +1194,14 @@ static int targetDataContiguous(ident_t *Loc, DeviceTy &Device, void *ArgsBase,
}
if (TPR.Flags.IsHostPointer) {
- DP("hst data:" DPxMOD " unified and shared, becomes a noop\n",
- DPxPTR(HstPtrBegin));
+ ODBG(ODT_Mapping) << "hst data:" << HstPtrBegin
+ << " unified and shared, becomes a noop";
return OFFLOAD_SUCCESS;
}
if (ArgType & OMP_TGT_MAPTYPE_TO) {
- DP("Moving %" PRId64 " bytes (hst:" DPxMOD ") -> (tgt:" DPxMOD ")\n",
- ArgSize, DPxPTR(HstPtrBegin), DPxPTR(TgtPtrBegin));
+ ODBG(ODT_Mapping) << "Moving " << ArgSize << " bytes (hst:" << HstPtrBegin
+ << ") -> (tgt:" << TgtPtrBegin << ")";
int Ret = Device.submitData(TgtPtrBegin, HstPtrBegin, ArgSize, AsyncInfo,
TPR.getEntry());
if (Ret != OFFLOAD_SUCCESS) {
@@ -1197,16 +1213,16 @@ static int targetDataContiguous(ident_t *Loc, DeviceTy &Device, void *ArgsBase,
[&](ShadowPtrInfoTy &ShadowPtr) {
constexpr int64_t VoidPtrSize = sizeof(void *);
if (ShadowPtr.PtrSize > VoidPtrSize) {
- DP("Restoring target descriptor " DPxMOD
- " to its original content (%" PRId64
- " bytes), containing pointee address " DPxMOD "\n",
- DPxPTR(ShadowPtr.TgtPtrAddr), ShadowPtr.PtrSize,
- DPxPTR(ShadowPtr.TgtPtrContent.data()));
+ ODBG(ODT_Mapping)
+ << "Restoring target descriptor " << ShadowPtr.TgtPtrAddr
+ << " to its original content (" << ShadowPtr.PtrSize
+ << " bytes), containing pointee address "
+ << ShadowPtr.TgtPtrContent.data();
} else {
- DP("Restoring target pointer " DPxMOD
- " to its original value " DPxMOD "\n",
- DPxPTR(ShadowPtr.TgtPtrAddr),
- DPxPTR(ShadowPtr.TgtPtrContent.data()));
+ ODBG(ODT_Mapping)
+ << "Restoring target pointer " << ShadowPtr.TgtPtrAddr
+ << " to its original value "
+ << ShadowPtr.TgtPtrContent.data();
}
Ret = Device.submitData(ShadowPtr.TgtPtrAddr,
ShadowPtr.TgtPtrContent.data(),
@@ -1218,15 +1234,15 @@ static int targetDataContiguous(ident_t *Loc, DeviceTy &Device, void *ArgsBase,
return OFFLOAD_SUCCESS;
});
if (Ret != OFFLOAD_SUCCESS) {
- DP("Updating shadow map failed\n");
+ ODBG(ODT_Mapping) << "Updating shadow map failed";
return Ret;
}
}
}
if (ArgType & OMP_TGT_MAPTYPE_FROM) {
- DP("Moving %" PRId64 " bytes (tgt:" DPxMOD ") -> (hst:" DPxMOD ")\n",
- ArgSize, DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin));
+ ODBG(ODT_Mapping) << "Moving " << ArgSize << " bytes (tgt:" << TgtPtrBegin
+ << ") -> (hst:" << HstPtrBegin << ")";
int Ret = Device.retrieveData(HstPtrBegin, TgtPtrBegin, ArgSize, AsyncInfo,
TPR.getEntry());
if (Ret != OFFLOAD_SUCCESS) {
@@ -1242,16 +1258,16 @@ static int targetDataContiguous(ident_t *Loc, DeviceTy &Device, void *ArgsBase,
[&](const ShadowPtrInfoTy &ShadowPtr) {
constexpr int64_t VoidPtrSize = sizeof(void *);
if (ShadowPtr.PtrSize > VoidPtrSize) {
- DP("Restoring host descriptor " DPxMOD
- " to its original content (%" PRId64
- " bytes), containing pointee address " DPxMOD "\n",
- DPxPTR(ShadowPtr.HstPtrAddr), ShadowPtr.PtrSize,
- DPxPTR(ShadowPtr.HstPtrContent.data()));
+ ODBG(ODT_Mapping)
+ << "Restoring host descriptor " << ShadowPtr.HstPtrAddr
+ << " to its original content (" << ShadowPtr.PtrSize
+ << " bytes), containing pointee address "
+ << ShadowPtr.HstPtrContent.data();
} else {
- DP("Restoring host pointer " DPxMOD
- " to its original value " DPxMOD "\n",
- DPxPTR(ShadowPtr.HstPtrAddr),
- DPxPTR(ShadowPtr.HstPtrContent.data()));
+ ODBG(ODT_Mapping)
+ << "Restoring host pointer " << ShadowPtr.HstPtrAddr
+ << " to its original value "
+ << ShadowPtr.HstPtrContent.data();
}
std::memcpy(ShadowPtr.HstPtrAddr, ShadowPtr.HstPtrContent.data(),
ShadowPtr.PtrSize);
@@ -1259,7 +1275,7 @@ static int targetDataContiguous(ident_t *Loc, DeviceTy &Device, void *ArgsBase,
});
Entry->unlock();
if (Ret != OFFLOAD_SUCCESS) {
- DP("Updating shadow map failed\n");
+ ODBG(ODT_Mapping) << "Updating shadow map failed";
return Ret;
}
return OFFLOAD_SUCCESS;
@@ -1295,9 +1311,8 @@ static int targetDataNonContiguous(ident_t *Loc, DeviceTy &Device,
}
} else {
char *Ptr = (char *)ArgsBase + Offset;
- DP("Transfer of non-contiguous : host ptr " DPxMOD " offset %" PRIu64
- " len %" PRIu64 "\n",
- DPxPTR(Ptr), Offset, Size);
+ ODBG(ODT_Mapping) << "Transfer of non-contiguous : host ptr " << Ptr
+ << " offset " << Offset << " len " << Size;
Ret = targetDataContiguous(Loc, Device, ArgsBase, Ptr, Size, ArgType,
AsyncInfo);
}
@@ -1330,8 +1345,8 @@ int targetDataUpdate(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
// Instead of executing the regular path of targetDataUpdate, call the
// targetDataMapper variant which will call targetDataUpdate again
// with new arguments.
- DP("Calling targetDataMapper for the %dth argument\n", I);
-
+ ODBG(ODT_Mapping) << "Calling targetDataMapper for the " << I
+ << "th argument";
map_var_info_t ArgName = (!ArgNames) ? nullptr : ArgNames[I];
int Ret = targetDataMapper(Loc, Device, ArgsBase[I], Args[I], ArgSizes[I],
ArgTypes[I], ArgName, ArgMappers[I], AsyncInfo,
@@ -1474,8 +1489,9 @@ class PrivateArgumentManagerTy {
// See if the pointee's begin address has corresponding storage on device.
void *TgtPteeBegin = [&]() -> void * {
if (!HstPteeBegin) {
- DP("Corresponding-pointer-initialization: pointee begin address is "
- "null\n");
+ ODBG(ODT_Mapping)
+ << "Corresponding-pointer-initialization: pointee begin address is "
+ "null";
return nullptr;
}
@@ -1586,9 +1602,10 @@ class PrivateArgumentManagerTy {
HstPteeBegin);
// Store the target pointee base address to the first VoidPtrSize bytes
- DP("Initializing corresponding-pointer-initialization source buffer "
- "for " DPxMOD ", with pointee base " DPxMOD "\n",
- DPxPTR(HstPtr), DPxPTR(TgtPteeBase));
+ ODBG(ODT_Mapping)
+ << "Corresponding-pointer-initialization: setting target pointee base "
+ "for "
+ << HstPtr << ", with pointee base " << TgtPteeBase;
std::memcpy(Buffer, &TgtPteeBase, VoidPtrSize);
if (HstPtrSize <= VoidPtrSize)
return;
@@ -1596,10 +1613,10 @@ class PrivateArgumentManagerTy {
// For Fortran descriptors, copy the remaining descriptor fields from host
uint64_t HstDescriptorFieldsSize = HstPtrSize - VoidPtrSize;
void *HstDescriptorFieldsAddr = static_cast<char *>(HstPtr) + VoidPtrSize;
- DP("Copying %" PRId64
- " bytes of descriptor fields into corresponding-pointer-initialization "
- "buffer at offset %" PRId64 ", from " DPxMOD "\n",
- HstDescriptorFieldsSize, VoidPtrSize, DPxPTR(HstDescriptorFieldsAddr));
+ ODBG(ODT_Mapping) << "Corresponding-pointer-initialization: copying "
+ << HstDescriptorFieldsSize
+ << " bytes of descriptor fields into buffer at offset "
+ << VoidPtrSize << ", from " << HstDescriptorFieldsAddr;
std::memcpy(Buffer + VoidPtrSize, HstDescriptorFieldsAddr,
HstDescriptorFieldsSize);
}
@@ -1638,21 +1655,21 @@ class PrivateArgumentManagerTy {
AllocImmediately) {
TgtPtr = Device.allocData(ArgSize, HstPtr);
if (!TgtPtr) {
- DP("Data allocation for %sprivate array " DPxMOD " failed.\n",
- (IsFirstPrivate ? "first-" : ""), DPxPTR(HstPtr));
+ ODBG(ODT_Alloc) << "Data allocation for "
+ << (IsFirstPrivate ? "first-" : "") << "private array "
+ << HstPtr << " failed.";
return OFFLOAD_FAIL;
}
-#ifdef OMPTARGET_DEBUG
- void *TgtPtrBase = (void *)((intptr_t)TgtPtr + ArgOffset);
- DP("Allocated %" PRId64 " bytes of target memory at " DPxMOD
- " for %sprivate array " DPxMOD " - pushing target argument " DPxMOD
- "\n",
- ArgSize, DPxPTR(TgtPtr), (IsFirstPrivate ? "first-" : ""),
- DPxPTR(HstPtr), DPxPTR(TgtPtrBase));
-#endif
+
+ ODBG(ODT_Alloc) << "Allocated " << ArgSize
+ << " bytes of target memory at " << TgtPtr << " for "
+ << (IsFirstPrivate ? "first-" : "") << "private array "
+ << HstPtr << " - pushing target argument "
+ << (void *)((intptr_t)TgtPtr + ArgOffset);
+
// If first-private, copy data from host
if (IsFirstPrivate) {
- DP("Submitting firstprivate data to the device.\n");
+ ODBG(ODT_Mapping) << "Submitting firstprivate data to the device.";
// The source value used for corresponding-pointer-initialization
// is different vs regular firstprivates.
@@ -1663,16 +1680,18 @@ class PrivateArgumentManagerTy {
: HstPtr;
int Ret = Device.submitData(TgtPtr, DataSource, ArgSize, AsyncInfo);
if (Ret != OFFLOAD_SUCCESS) {
- DP("Copying %s data to device failed.\n",
- IsCorrespondingPointerInit ? "corresponding-pointer-initialization"
- : "firstprivate");
+ ODBG(ODT_Mapping) << "Copying "
+ << (IsCorrespondingPointerInit
+ ? "corresponding-pointer-initialization"
+ : "firstprivate")
+ << " data to device failed.";
return OFFLOAD_FAIL;
}
}
TgtPtrs.push_back(TgtPtr);
} else {
- DP("Firstprivate array " DPxMOD " of size %" PRId64 " will be packed\n",
- DPxPTR(HstPtr), ArgSize);
+ ODBG(ODT_Mapping) << "Firstprivate array " << HstPtr << " of size "
+ << ArgSize << " will be packed";
// When reach this point, the argument must meet all following
// requirements:
// 1. Its size does not exceed the threshold (see the comment for
@@ -1746,17 +1765,18 @@ class PrivateArgumentManagerTy {
void *TgtPtr =
Device.allocData(FirstPrivateArgSize, FirstPrivateArgBuffer.data());
if (TgtPtr == nullptr) {
- DP("Failed to allocate target memory for private arguments.\n");
+ ODBG(ODT_Alloc)
+ << "Failed to allocate target memory for private arguments.";
return OFFLOAD_FAIL;
}
TgtPtrs.push_back(TgtPtr);
- DP("Allocated %" PRId64 " bytes of target memory at " DPxMOD "\n",
- FirstPrivateArgSize, DPxPTR(TgtPtr));
+ ODBG(ODT_Alloc) << "Allocated " << FirstPrivateArgSize
+ << " bytes of target memory at " << TgtPtr;
// Transfer data to target device
int Ret = Device.submitData(TgtPtr, FirstPrivateArgBuffer.data(),
FirstPrivateArgSize, AsyncInfo);
if (Ret != OFFLOAD_SUCCESS) {
- DP("Failed to submit data of private arguments.\n");
+ ODBG(ODT_DataTransfer) << "Failed to submit data of private arguments.";
return OFFLOAD_FAIL;
}
// Fill in all placeholder pointers
@@ -1768,10 +1788,9 @@ class PrivateArgumentManagerTy {
TP += Info.Padding;
Ptr = reinterpret_cast<void *>(TP);
TP += Info.Size;
- DP("Firstprivate array " DPxMOD " of size %" PRId64 " mapped to " DPxMOD
- "\n",
- DPxPTR(Info.HstPtrBegin), Info.HstPtrEnd - Info.HstPtrBegin,
- DPxPTR(Ptr));
+ ODBG(ODT_Mapping) << "Firstprivate array " << Info.HstPtrBegin
+ << " of size " << (Info.HstPtrEnd - Info.HstPtrBegin)
+ << " mapped to " << Ptr;
}
}
@@ -1783,7 +1802,7 @@ class PrivateArgumentManagerTy {
for (void *P : TgtPtrs) {
int Ret = Device.deleteData(P);
if (Ret != OFFLOAD_SUCCESS) {
- DP("Deallocation of (first-)private arrays failed.\n");
+ ODBG(ODT_Alloc) << "Deallocation of (first-)private arrays failed.";
return OFFLOAD_FAIL;
}
}
@@ -1851,7 +1870,7 @@ static int processDataBefore(ident_t *Loc, int64_t DeviceId, void *HostPtr,
void *HstPtrBase = Args[Idx];
void *TgtPtrBase =
(void *)((intptr_t)TgtArgs[TgtIdx] + TgtOffsets[TgtIdx]);
- DP("Parent lambda base " DPxMOD "\n", DPxPTR(TgtPtrBase));
+ ODBG(ODT_Mapping) << "Parent lambda base " << TgtPtrBase;
uint64_t Delta = (uint64_t)HstPtrBegin - (uint64_t)HstPtrBase;
void *TgtPtrBegin = (void *)((uintptr_t)TgtPtrBase + Delta);
void *&PointerTgtPtrBegin = AsyncInfo.getVoidPtrLocation();
@@ -1861,18 +1880,19 @@ static int processDataBefore(ident_t *Loc, int64_t DeviceId, void *HostPtr,
/*UseHoldRefCount=*/false);
PointerTgtPtrBegin = TPR.TargetPointer;
if (!TPR.isPresent()) {
- DP("No lambda captured variable mapped (" DPxMOD ") - ignored\n",
- DPxPTR(HstPtrVal));
+ ODBG(ODT_Mapping) << "No lambda captured variable mapped "
+ << HstPtrVal << " - ignored";
continue;
}
if (TPR.Flags.IsHostPointer) {
- DP("Unified memory is active, no need to map lambda captured"
- "variable (" DPxMOD ")\n",
- DPxPTR(HstPtrVal));
+ ODBG(ODT_Mapping)
+ << "Unified memory is active, no need to map lambda captured"
+ "variable ("
+ << HstPtrVal << ")";
continue;
}
- DP("Update lambda reference (" DPxMOD ") -> [" DPxMOD "]\n",
- DPxPTR(PointerTgtPtrBegin), DPxPTR(TgtPtrBegin));
+ ODBG(ODT_Mapping) << "Update lambda reference (" << PointerTgtPtrBegin
+ << ") -> [" << TgtPtrBegin << "]";
Ret =
DeviceOrErr->submitData(TgtPtrBegin, &PointerTgtPtrBegin,
sizeof(void *), AsyncInfo, TPR.getEntry());
@@ -1890,8 +1910,8 @@ static int processDataBefore(ident_t *Loc, int64_t DeviceId, void *HostPtr,
ptrdiff_t TgtBaseOffset;
TargetPointerResultTy TPR;
if (ArgTypes[I] & OMP_TGT_MAPTYPE_LITERAL) {
- DP("Forwarding first-private value " DPxMOD " to the target construct\n",
- DPxPTR(HstPtrBase));
+ ODBG(ODT_Mapping) << "Forwarding first-private value " << HstPtrBase
+ << " to the target construct";
TgtPtrBegin = HstPtrBase;
TgtBaseOffset = 0;
} else if (ArgTypes[I] & OMP_TGT_MAPTYPE_PRIVATE) {
@@ -1955,11 +1975,9 @@ static int processDataBefore(ident_t *Loc, int64_t DeviceId, void *HostPtr,
/*UseHoldRefCount=*/false);
TgtPtrBegin = TPR.TargetPointer;
TgtBaseOffset = (intptr_t)HstPtrBase - (intptr_t)HstPtrBegin;
-#ifdef OMPTARGET_DEBUG
void *TgtPtrBase = (void *)((intptr_t)TgtPtrBegin + TgtBaseOffset);
- DP("Obtained target argument " DPxMOD " from host pointer " DPxMOD "\n",
- DPxPTR(TgtPtrBase), DPxPTR(HstPtrBegin));
-#endif
+ ODBG(ODT_Mapping) << "Obtained target argument " << TgtPtrBase
+ << " from host pointer " << HstPtrBegin;
}
TgtArgsPositions[I] = TgtArgs.size();
TgtArgs.push_back(TgtPtrBegin);
@@ -1972,7 +1990,7 @@ static int processDataBefore(ident_t *Loc, int64_t DeviceId, void *HostPtr,
// Pack and transfer first-private arguments
Ret = PrivateArgumentManager.packAndTransfer(TgtArgs);
if (Ret != OFFLOAD_SUCCESS) {
- DP("Failed to pack and transfer first private arguments\n");
+ ODBG(ODT_Mapping) << "Failed to pack and transfer first private arguments";
return OFFLOAD_FAIL;
}
@@ -2030,8 +2048,8 @@ int target(ident_t *Loc, DeviceTy &Device, void *HostPtr,
TableMap *TM = getTableMap(HostPtr);
// No map for this host pointer found!
if (!TM) {
- REPORT("Host ptr " DPxMOD " does not have a matching target pointer.\n",
- DPxPTR(HostPtr));
+ REPORT() << "Host ptr " << HostPtr
+ << " does not have a matching target pointer.";
return OFFLOAD_FAIL;
}
@@ -2045,7 +2063,7 @@ int target(ident_t *Loc, DeviceTy &Device, void *HostPtr,
}
assert(TargetTable && "Global data has not been mapped\n");
- DP("loop trip count is %" PRIu64 ".\n", KernelArgs.Tripcount);
+ ODBG(ODT_Kernel) << "loop trip count is " << KernelArgs.Tripcount;
// We need to keep bases and offsets separate. Sometimes (e.g. in OpenCL) we
// need to manifest base pointers prior to launching a kernel. Even if we have
@@ -2084,9 +2102,10 @@ int target(ident_t *Loc, DeviceTy &Device, void *HostPtr,
// Launch device execution.
void *TgtEntryPtr = TargetTable->EntriesBegin[TM->Index].Address;
- DP("Launching target execution %s with pointer " DPxMOD " (index=%d).\n",
- TargetTable->EntriesBegin[TM->Index].SymbolName, DPxPTR(TgtEntryPtr),
- TM->Index);
+ ODBG(ODT_Kernel) << "Launching target execution "
+ << TargetTable->EntriesBegin[TM->Index].SymbolName
+ << " with pointer " << TgtEntryPtr << " (index=" << TM->Index
+ << ").";
{
assert(KernelArgs.NumArgs == TgtArgs.size() && "Argument count mismatch!");
@@ -2110,7 +2129,7 @@ int target(ident_t *Loc, DeviceTy &Device, void *HostPtr,
}
if (Ret != OFFLOAD_SUCCESS) {
- REPORT("Executing target region abort target.\n");
+ REPORT() << "Executing target region abort target.";
return OFFLOAD_FAIL;
}
@@ -2155,8 +2174,8 @@ int target_replay(ident_t *Loc, DeviceTy &Device, void *HostPtr,
// Fail if the table map fails to find the target kernel pointer for the
// provided host pointer.
if (!TM) {
- REPORT("Host ptr " DPxMOD " does not have a matching target pointer.\n",
- DPxPTR(HostPtr));
+ REPORT() << "Host ptr " << HostPtr
+ << " does not have a matching target pointer.";
return OFFLOAD_FAIL;
}
@@ -2173,9 +2192,10 @@ int target_replay(ident_t *Loc, DeviceTy &Device, void *HostPtr,
// Retrieve the target kernel pointer, allocate and store the recorded device
// memory data, and launch device execution.
void *TgtEntryPtr = TargetTable->EntriesBegin[TM->Index].Address;
- DP("Launching target execution %s with pointer " DPxMOD " (index=%d).\n",
- TargetTable->EntriesBegin[TM->Index].SymbolName, DPxPTR(TgtEntryPtr),
- TM->Index);
+ ODBG(ODT_Kernel) << "Launching target execution "
+ << TargetTable->EntriesBegin[TM->Index].SymbolName
+ << " with pointer " << TgtEntryPtr << " (index=" << TM->Index
+ << ").";
void *TgtPtr = Device.allocData(DeviceMemorySize, /*HstPtr=*/nullptr,
TARGET_ALLOC_DEFAULT);
@@ -2192,7 +2212,7 @@ int target_replay(ident_t *Loc, DeviceTy &Device, void *HostPtr,
AsyncInfo);
if (Ret != OFFLOAD_SUCCESS) {
- REPORT("Executing target region abort target.\n");
+ REPORT() << "Executing target region abort target.";
return OFFLOAD_FAIL;
}
>From ce0bfc8055a4afb801b36db3848ad7d3bc5e14ea Mon Sep 17 00:00:00 2001
From: Alex Duran <alejandro.duran at intel.com>
Date: Wed, 3 Dec 2025 15:40:45 +0100
Subject: [PATCH 6/7] remove leftover function
---
offload/include/Shared/Debug.h | 2 --
1 file changed, 2 deletions(-)
diff --git a/offload/include/Shared/Debug.h b/offload/include/Shared/Debug.h
index d60ba26257bbe..bda611abdb46c 100644
--- a/offload/include/Shared/Debug.h
+++ b/offload/include/Shared/Debug.h
@@ -601,6 +601,4 @@ static inline odbg_ostream reportErrorStream() {
} // namespace llvm::omp::target::debug
-inline int getDebugLevel() { return 1; }
-
#endif // OMPTARGET_SHARED_DEBUG_H
>From 8f5eb8d3431b2de7082eeb39daf79fe0958f4645 Mon Sep 17 00:00:00 2001
From: Alex Duran <alejandro.duran at intel.com>
Date: Wed, 10 Dec 2025 14:13:19 +0100
Subject: [PATCH 7/7] refactor LambdaHelper
---
offload/include/Shared/Debug.h | 27 +++++++++------------------
1 file changed, 9 insertions(+), 18 deletions(-)
diff --git a/offload/include/Shared/Debug.h b/offload/include/Shared/Debug.h
index bda611abdb46c..c30c2106bca46 100644
--- a/offload/include/Shared/Debug.h
+++ b/offload/include/Shared/Debug.h
@@ -431,27 +431,18 @@ static inline raw_ostream &operator<<(raw_ostream &Os,
static_cast<llvm::offload::debug::odbg_ostream::IfLevel>(0)
// helper templates to support lambdas with different number of arguments
-
-template <typename LambdaTy> struct lambdaHelper {
- template <typename FuncTy, typename RetTy, typename... Args>
- static constexpr size_t CountArgs(RetTy (FuncTy::*)(Args...)) {
- return sizeof...(Args);
- }
-
- template <typename FuncTy, typename RetTy, typename... Args>
- static constexpr size_t CountArgs(RetTy (FuncTy::*)(Args...) const) {
- return sizeof...(Args);
- }
-
- static constexpr size_t NArgs = CountArgs(&LambdaTy::operator());
+template <typename LambdaTy> struct LambdaHelper {
+ template <typename T, typename = std::void_t<>>
+ struct has_two_args : std::false_type {};
+ template <typename T>
+ struct has_two_args<T, std::void_t<decltype(std::declval<T>().operator()(1,2))>>
+ : std::true_type {};
static void dispatch(LambdaTy func, llvm::raw_ostream &Os, uint32_t Level) {
- if constexpr (NArgs == 1)
- func(Os);
- else if constexpr (NArgs == 2)
+ if constexpr (has_two_args<LambdaTy>::value)
func(Os, Level);
else
- static_assert(true, "Unsupported number of arguments in debug callback");
+ func(Os);
}
};
@@ -465,7 +456,7 @@ template <typename LambdaTy> struct lambdaHelper {
RealLevel, /*ShouldPrefixNextString=*/true, \
/*ShouldEmitNewLineOnDestruction=*/true}; \
auto F = Callback; \
- ::llvm::offload::debug::lambdaHelper<decltype(F)>::dispatch(F, OS, \
+ ::llvm::offload::debug::LambdaHelper<decltype(F)>::dispatch(F, OS, \
RealLevel); \
} \
}
More information about the llvm-commits
mailing list