From f63c8e6027290d797a99a00c6574f08d68cb09b1 Mon Sep 17 00:00:00 2001 From: Alex Duran Date: Tue, 2 Dec 2025 09:56:37 +0100 Subject: [PATCH 1/8] [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 +#include #include #include @@ -78,17 +79,6 @@ inline std::atomic &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 - -#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 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 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 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 Date: Tue, 2 Dec 2025 10:05:01 +0100 Subject: [PATCH 2/8] 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 Date: Tue, 2 Dec 2025 10:07:44 +0100 Subject: [PATCH 3/8] 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 Date: Wed, 3 Dec 2025 06:16:18 +0100 Subject: [PATCH 4/8] [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(0) +// helper templates to support lambdas with different number of arguments + +template struct lambdaHelper { + template + static constexpr size_t CountArgs(RetTy (FuncTy::*)(Args...)) { + return sizeof...(Args); + } + + template + 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::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 #include @@ -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, "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( reinterpret_cast(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 Date: Wed, 3 Dec 2025 06:52:00 +0100 Subject: [PATCH 5/8] 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(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( reinterpret_cast(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(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(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 Date: Wed, 3 Dec 2025 15:40:45 +0100 Subject: [PATCH 6/8] 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 Date: Wed, 10 Dec 2025 14:13:19 +0100 Subject: [PATCH 7/8] 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(0) // helper templates to support lambdas with different number of arguments - -template struct lambdaHelper { - template - static constexpr size_t CountArgs(RetTy (FuncTy::*)(Args...)) { - return sizeof...(Args); - } - - template - static constexpr size_t CountArgs(RetTy (FuncTy::*)(Args...) const) { - return sizeof...(Args); - } - - static constexpr size_t NArgs = CountArgs(&LambdaTy::operator()); +template struct LambdaHelper { + template > + struct has_two_args : std::false_type {}; + template + struct has_two_args().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::value) func(Os, Level); else - static_assert(true, "Unsupported number of arguments in debug callback"); + func(Os); } }; @@ -465,7 +456,7 @@ template struct lambdaHelper { RealLevel, /*ShouldPrefixNextString=*/true, \ /*ShouldEmitNewLineOnDestruction=*/true}; \ auto F = Callback; \ - ::llvm::offload::debug::lambdaHelper::dispatch(F, OS, \ + ::llvm::offload::debug::LambdaHelper::dispatch(F, OS, \ RealLevel); \ } \ } From f15c64c7aed2a122425864a883ed7b5c8db7a6f2 Mon Sep 17 00:00:00 2001 From: Alex Duran Date: Wed, 10 Dec 2025 14:15:42 +0100 Subject: [PATCH 8/8] format --- offload/include/Shared/Debug.h | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/offload/include/Shared/Debug.h b/offload/include/Shared/Debug.h index c30c2106bca46..d5cf719f1ebf3 100644 --- a/offload/include/Shared/Debug.h +++ b/offload/include/Shared/Debug.h @@ -435,7 +435,8 @@ template struct LambdaHelper { template > struct has_two_args : std::false_type {}; template - struct has_two_args().operator()(1,2))>> + struct has_two_args().operator()(1, 2))>> : std::true_type {}; static void dispatch(LambdaTy func, llvm::raw_ostream &Os, uint32_t Level) {