[llvm] [Offload][OpenMP] Report source location of data transfers and synchr… (PR #223730)
Jason Van Beusekom via llvm-commits
llvm-commits at lists.llvm.org
Tue Sep 29 13:53:15 PDT 2026
https://github.com/Jason-Van-Beusekom updated https://github.com/llvm/llvm-project/pull/223730
>From 1a6608d3583e001f27960d946781b07110c4e3ab Mon Sep 17 00:00:00 2001
From: Jason Van Beusekom <jason.van-beusekom at hpe.com>
Date: Tue, 15 Sep 2026 09:54:28 -0500
Subject: [PATCH 1/3] [Offload][OpenMP] Report source location of data
transfers and synchronization
---
offload/include/OpenMP/Mapping.h | 6 ++++--
offload/include/Shared/SourceInfo.h | 13 +++++++++++++
offload/include/device.h | 6 ++++--
offload/libompaccsupport/Mapping.cpp | 27 +++++++++++++++++----------
offload/libompaccsupport/device.cpp | 17 +++++++++++------
offload/libomptarget/interface.cpp | 8 ++++++--
offload/libomptarget/omptarget.cpp | 9 +++++----
offload/libomptarget/private.h | 11 +++++++++++
offload/test/offloading/force-usm.cpp | 1 +
offload/test/offloading/info.c | 10 +++++++---
10 files changed, 79 insertions(+), 29 deletions(-)
diff --git a/offload/include/OpenMP/Mapping.h b/offload/include/OpenMP/Mapping.h
index e4024abf26690..4b160684ddb32 100644
--- a/offload/include/OpenMP/Mapping.h
+++ b/offload/include/OpenMP/Mapping.h
@@ -671,7 +671,8 @@ struct MappingInfoTy {
bool HasFlagTo, bool HasFlagAlways, bool IsImplicit, bool UpdateRefCount,
bool HasCloseModifier, bool HasPresentModifier, bool HasHoldModifier,
AsyncInfoTy &AsyncInfo, HostDataToTargetTy *OwnedTPR = nullptr,
- bool ReleaseHDTTMap = true, StateInfoTy *StateInfo = nullptr);
+ bool ReleaseHDTTMap = true, StateInfoTy *StateInfo = nullptr,
+ const ident_t *Loc = nullptr);
/// Return the target pointer for \p HstPtrBegin in \p HDTTMap. The accessor
/// ensures exclusive access to the HDTT map.
@@ -717,7 +718,8 @@ struct MappingInfoTy {
/// is set, the associated metadata will be printed as well.
void printCopyInfo(void *TgtPtr, void *HstPtr, int64_t Size, bool H2D,
HostDataToTargetTy *Entry,
- MappingInfoTy::HDTTMapAccessorTy *HDTTMapPtr);
+ MappingInfoTy::HDTTMapAccessorTy *HDTTMapPtr,
+ const ident_t *Loc = nullptr, const char *Name = nullptr);
private:
DeviceTy &Device;
diff --git a/offload/include/Shared/SourceInfo.h b/offload/include/Shared/SourceInfo.h
index e8dfb46c973e0..2b42dec37549f 100644
--- a/offload/include/Shared/SourceInfo.h
+++ b/offload/include/Shared/SourceInfo.h
@@ -99,4 +99,17 @@ static inline std::string getNameFromMapping(const map_var_info_t Name) {
return NameStr.substr(Begin + 1, End - Begin - 1);
}
+/// Returns "<Prefix><filename>:<line>:<column>" when \p Loc carries a
+/// compiler-provided source location, or an empty string otherwise, so info
+/// output can be annotated without emitting "unknown:0:0".
+static inline std::string getSourceLocationSuffix(const ident_t *Loc,
+ const char *Prefix) {
+ SourceInfo Info(Loc);
+ if (!Info.isAvailible())
+ return "";
+ return std::string(Prefix) + Info.getFilename() + ":" +
+ std::to_string(Info.getLine()) + ":" +
+ std::to_string(Info.getColumn());
+}
+
#endif // OMPTARGET_SHARED_SOURCE_INFO_H
diff --git a/offload/include/device.h b/offload/include/device.h
index 5918d04d9e0d4..c90e0eaf7f2cb 100644
--- a/offload/include/device.h
+++ b/offload/include/device.h
@@ -87,13 +87,15 @@ struct DeviceTy {
int32_t submitData(void *TgtPtrBegin, void *HstPtrBegin, int64_t Size,
AsyncInfoTy &AsyncInfo,
HostDataToTargetTy *Entry = nullptr,
- MappingInfoTy::HDTTMapAccessorTy *HDTTMapPtr = nullptr);
+ MappingInfoTy::HDTTMapAccessorTy *HDTTMapPtr = nullptr,
+ const ident_t *Loc = nullptr, const char *Name = nullptr);
// Copy data from device back to host
int32_t retrieveData(void *HstPtrBegin, void *TgtPtrBegin, int64_t Size,
AsyncInfoTy &AsyncInfo,
HostDataToTargetTy *Entry = nullptr,
- MappingInfoTy::HDTTMapAccessorTy *HDTTMapPtr = nullptr);
+ MappingInfoTy::HDTTMapAccessorTy *HDTTMapPtr = nullptr,
+ const ident_t *Loc = nullptr);
// Return true if data can be copied to DstDevice directly
bool isDataExchangable(const DeviceTy &DstDevice);
diff --git a/offload/libompaccsupport/Mapping.cpp b/offload/libompaccsupport/Mapping.cpp
index 1bb2e424bd083..a4905e4106deb 100644
--- a/offload/libompaccsupport/Mapping.cpp
+++ b/offload/libompaccsupport/Mapping.cpp
@@ -210,7 +210,7 @@ TargetPointerResultTy MappingInfoTy::getTargetPointer(
bool HasFlagAlways, bool IsImplicit, bool UpdateRefCount,
bool HasCloseModifier, bool HasPresentModifier, bool HasHoldModifier,
AsyncInfoTy &AsyncInfo, HostDataToTargetTy *OwnedTPR, bool ReleaseHDTTMap,
- StateInfoTy *StateInfo) {
+ StateInfoTy *StateInfo, const ident_t *Loc) {
LookupResult LR = lookupMapping(HDTTMap, HstPtrBegin, Size, OwnedTPR);
LR.TPR.Flags.IsPresent = true;
@@ -385,7 +385,8 @@ TargetPointerResultTy MappingInfoTy::getTargetPointer(
<< ") -> (tgt:" << LR.TPR.TargetPointer << ")";
int Ret = Device.submitData(LR.TPR.TargetPointer, HstPtrBegin, Size,
- AsyncInfo, LR.TPR.getEntry());
+ AsyncInfo, LR.TPR.getEntry(),
+ /*HDTTMapPtr=*/nullptr, Loc);
if (Ret != OFFLOAD_SUCCESS) {
REPORT() << "Copying data to device failed.";
// We will also return nullptr if the data movement fails because that
@@ -557,21 +558,27 @@ int MappingInfoTy::deallocTgtPtrAndEntry(HostDataToTargetTy *Entry,
static void printCopyInfoImpl(int DeviceId, bool H2D, void *SrcPtrBegin,
void *DstPtrBegin, int64_t Size,
- HostDataToTargetTy *HT) {
+ HostDataToTargetTy *HT, const ident_t *Loc,
+ const char *Name) {
+
+ std::string LocStr = getSourceLocationSuffix(Loc, ", at ");
INFO(OMP_INFOTYPE_DATA_TRANSFER, DeviceId,
"Copying data from %s to %s, %sPtr=" DPxMOD ", %sPtr=" DPxMOD
- ", Size=%" PRId64 ", Name=%s\n",
+ ", Size=%" PRId64 ", Name=%s%s\n",
H2D ? "host" : "device", H2D ? "device" : "host", H2D ? "Hst" : "Tgt",
DPxPTR(H2D ? SrcPtrBegin : DstPtrBegin), H2D ? "Tgt" : "Hst",
DPxPTR(H2D ? DstPtrBegin : SrcPtrBegin), Size,
(HT && HT->HstPtrName) ? getNameFromMapping(HT->HstPtrName).c_str()
- : "unknown");
+ : (Name ? Name : "unknown"),
+ LocStr.c_str());
}
-void MappingInfoTy::printCopyInfo(
- void *TgtPtrBegin, void *HstPtrBegin, int64_t Size, bool H2D,
- HostDataToTargetTy *Entry, MappingInfoTy::HDTTMapAccessorTy *HDTTMapPtr) {
+void MappingInfoTy::printCopyInfo(void *TgtPtrBegin, void *HstPtrBegin,
+ int64_t Size, bool H2D,
+ HostDataToTargetTy *Entry,
+ MappingInfoTy::HDTTMapAccessorTy *HDTTMapPtr,
+ const ident_t *Loc, const char *Name) {
auto HDTTMap =
HostDataToTargetMap.getExclusiveAccessor(!!Entry || !!HDTTMapPtr);
LookupResult LR;
@@ -579,6 +586,6 @@ void MappingInfoTy::printCopyInfo(
LR = lookupMapping(HDTTMapPtr ? *HDTTMapPtr : HDTTMap, HstPtrBegin, Size);
Entry = LR.TPR.getEntry();
}
- printCopyInfoImpl(Device.DeviceID, H2D, HstPtrBegin, TgtPtrBegin, Size,
- Entry);
+ printCopyInfoImpl(Device.DeviceID, H2D, HstPtrBegin, TgtPtrBegin, Size, Entry,
+ Loc, Name);
}
diff --git a/offload/libompaccsupport/device.cpp b/offload/libompaccsupport/device.cpp
index 5aec1156930d6..3dd8106405adb 100644
--- a/offload/libompaccsupport/device.cpp
+++ b/offload/libompaccsupport/device.cpp
@@ -206,7 +206,8 @@ setupIndirectCallTable(DeviceTy &Device, __tgt_device_image *Image,
IndirectCallTable.size() * sizeof(std::pair<void *, void *>);
void *DevicePtr = Device.allocData(TableSize, nullptr, TARGET_ALLOC_DEVICE);
if (Device.submitData(DevicePtr, IndirectCallTable.data(), TableSize,
- AsyncInfo))
+ AsyncInfo, /*Entry=*/nullptr, /*HDTTMapPtr=*/nullptr,
+ /*Loc=*/nullptr, /*Name=*/"IndirectCallTable"))
return error::createOffloadError(error::ErrorCode::INVALID_BINARY,
"failed to copy data");
// The IndirectCallTable is on the stack, so we must synchronize to ensure
@@ -263,7 +264,9 @@ DeviceTy::loadBinary(__tgt_device_image *Img) {
AsyncInfoTy AsyncInfo(*this);
if (submitData(DeviceEnvironmentPtr, &DeviceEnvironment,
- sizeof(DeviceEnvironment), AsyncInfo))
+ sizeof(DeviceEnvironment), AsyncInfo, /*Entry=*/nullptr,
+ /*HDTTMapPtr=*/nullptr, /*Loc=*/nullptr,
+ /*Name=*/"DeviceEnvironment"))
return error::createOffloadError(error::ErrorCode::INVALID_BINARY,
"failed to copy data");
@@ -295,10 +298,11 @@ int32_t DeviceTy::deleteData(void *TgtAllocBegin, int32_t Kind) {
// Submit data to device
int32_t DeviceTy::submitData(void *TgtPtrBegin, void *HstPtrBegin, int64_t Size,
AsyncInfoTy &AsyncInfo, HostDataToTargetTy *Entry,
- MappingInfoTy::HDTTMapAccessorTy *HDTTMapPtr) {
+ MappingInfoTy::HDTTMapAccessorTy *HDTTMapPtr,
+ const ident_t *Loc, const char *Name) {
if (getInfoLevel() & OMP_INFOTYPE_DATA_TRANSFER)
MappingInfo.printCopyInfo(TgtPtrBegin, HstPtrBegin, Size, /*H2D=*/true,
- Entry, HDTTMapPtr);
+ Entry, HDTTMapPtr, Loc, Name);
/// RAII to establish tool anchors before and after data submit
OMPT_IF_BUILT(
@@ -315,10 +319,11 @@ int32_t DeviceTy::submitData(void *TgtPtrBegin, void *HstPtrBegin, int64_t Size,
int32_t DeviceTy::retrieveData(void *HstPtrBegin, void *TgtPtrBegin,
int64_t Size, AsyncInfoTy &AsyncInfo,
HostDataToTargetTy *Entry,
- MappingInfoTy::HDTTMapAccessorTy *HDTTMapPtr) {
+ MappingInfoTy::HDTTMapAccessorTy *HDTTMapPtr,
+ const ident_t *Loc) {
if (getInfoLevel() & OMP_INFOTYPE_DATA_TRANSFER)
MappingInfo.printCopyInfo(TgtPtrBegin, HstPtrBegin, Size, /*H2D=*/false,
- Entry, HDTTMapPtr);
+ Entry, HDTTMapPtr, Loc);
/// RAII to establish tool anchors before and after data retrieval
OMPT_IF_BUILT(
diff --git a/offload/libomptarget/interface.cpp b/offload/libomptarget/interface.cpp
index 5833925209fee..51e7b6109c094 100644
--- a/offload/libomptarget/interface.cpp
+++ b/offload/libomptarget/interface.cpp
@@ -186,8 +186,10 @@ targetData(ident_t *Loc, int64_t DeviceId, int32_t ArgNum, void **ArgsBase,
if (StateInfo && !StateInfo->AttachEntries.empty())
Rc = processAttachEntries(*DeviceOrErr, *StateInfo, AsyncInfo);
- if (Rc == OFFLOAD_SUCCESS)
+ if (Rc == OFFLOAD_SUCCESS) {
+ printSyncInfo(Loc, DeviceId);
Rc = AsyncInfo.synchronize();
+ }
}
handleTargetOutcome(Rc == OFFLOAD_SUCCESS, Loc);
@@ -439,8 +441,10 @@ static inline int targetKernel(ident_t *Loc, int64_t DeviceId, int32_t NumTeams,
Rc = target(Loc, *DeviceOrErr, HostPtr, *KernelArgs, AsyncInfo);
{ // required to show synchronization
TIMESCOPE_WITH_DETAILS_AND_IDENT("Runtime: synchronize", "", Loc);
- if (Rc == OFFLOAD_SUCCESS)
+ if (Rc == OFFLOAD_SUCCESS) {
+ printSyncInfo(Loc, DeviceId);
Rc = AsyncInfo.synchronize();
+ }
handleTargetOutcome(Rc == OFFLOAD_SUCCESS, Loc);
assert(Rc == OFFLOAD_SUCCESS && "__tgt_target_kernel unexpected failure!");
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index 2513ae8d19814..c2abbc95bf251 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -613,7 +613,7 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
HDTTMap, HstPtrBegin, HstPtrBase, TgtPadding, DataSize, HstPtrName,
HasFlagTo, HasFlagAlways, IsImplicit, UpdateRef, HasCloseModifier,
HasPresentModifier, HasHoldModifier, AsyncInfo, PointerTpr.getEntry(),
- /*ReleaseHDTTMap=*/true, StateInfo);
+ /*ReleaseHDTTMap=*/true, StateInfo, Loc);
void *TgtPtrBegin = TPR.TargetPointer;
IsHostPtr = TPR.Flags.IsHostPointer;
// If data_size==0, then the argument could be a zero-length pointer to
@@ -1184,7 +1184,8 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
}
}
- int Ret = Device.retrieveData(HstPtr, TgtPtr, Size, AsyncInfo, Entry);
+ int Ret = Device.retrieveData(HstPtr, TgtPtr, Size, AsyncInfo, Entry,
+ /*HDTTMapPtr=*/nullptr, Loc);
if (Ret != OFFLOAD_SUCCESS) {
REPORT() << "Copying data from device failed.";
return OFFLOAD_FAIL;
@@ -1365,7 +1366,7 @@ static int targetDataContiguous(ident_t *Loc, DeviceTy &Device, void *ArgsBase,
ODBG(ODT_Mapping) << "Moving " << ArgSize << " bytes (hst:" << HstPtrBegin
<< ") -> (tgt:" << TgtPtrBegin << ")";
int Ret = Device.submitData(TgtPtrBegin, HstPtrBegin, ArgSize, AsyncInfo,
- TPR.getEntry());
+ TPR.getEntry(), /*HDTTMapPtr=*/nullptr, Loc);
if (Ret != OFFLOAD_SUCCESS) {
REPORT() << "Copying data to device failed.";
return OFFLOAD_FAIL;
@@ -1406,7 +1407,7 @@ static int targetDataContiguous(ident_t *Loc, DeviceTy &Device, void *ArgsBase,
ODBG(ODT_Mapping) << "Moving " << ArgSize << " bytes (tgt:" << TgtPtrBegin
<< ") -> (hst:" << HstPtrBegin << ")";
int Ret = Device.retrieveData(HstPtrBegin, TgtPtrBegin, ArgSize, AsyncInfo,
- TPR.getEntry());
+ TPR.getEntry(), /*HDTTMapPtr=*/nullptr, Loc);
if (Ret != OFFLOAD_SUCCESS) {
REPORT() << "Copying data from device failed.";
return OFFLOAD_FAIL;
diff --git a/offload/libomptarget/private.h b/offload/libomptarget/private.h
index 80630f9c366ac..0de02cb2304d6 100644
--- a/offload/libomptarget/private.h
+++ b/offload/libomptarget/private.h
@@ -86,6 +86,17 @@ printKernelArguments(const ident_t *Loc, const int64_t DeviceId,
}
}
+////////////////////////////////////////////////////////////////////////////////
+/// Report that the runtime is about to wait for the region's outstanding
+/// asynchronous operations (data transfers and kernels) to complete.
+static inline void printSyncInfo(const ident_t *Loc, const int64_t DeviceId) {
+ if (!(getInfoLevel() & OMP_INFOTYPE_DATA_TRANSFER))
+ return;
+ std::string LocStr = getSourceLocationSuffix(Loc, " at ");
+ INFO(OMP_INFOTYPE_DATA_TRANSFER, DeviceId,
+ "Waiting for asynchronous operations to complete%s\n", LocStr.c_str());
+}
+
////////////////////////////////////////////////////////////////////////////////
/// Checks if the passed device is the initial device (i.e., host device)
/// While the device number is defined as the value of the total number of
diff --git a/offload/test/offloading/force-usm.cpp b/offload/test/offloading/force-usm.cpp
index 8a0fd1828b809..f8a2f1a80645e 100644
--- a/offload/test/offloading/force-usm.cpp
+++ b/offload/test/offloading/force-usm.cpp
@@ -55,6 +55,7 @@ int main(void) {
// NO-USM-NEXT: device 0 info: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}}, Size=4
// NO-USM-NEXT: device 0 info: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}}, Size=12
// NO-USM-NEXT: device 0 info: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}}, Size=4
+// NO-USM-NEXT: device 0 info: Waiting for asynchronous operations to complete
// NO-USM-NEXT: SUCCESS
// FORCE-USM: SUCCESS
diff --git a/offload/test/offloading/info.c b/offload/test/offloading/info.c
index 410662baf2b5b..cfe4ad3afaa10 100644
--- a/offload/test/offloading/info.c
+++ b/offload/test/offloading/info.c
@@ -30,11 +30,13 @@ int main() {
// INFO: info: alloc(A[0:64])[256]
// INFO: info: tofrom(B[0:64])[256]
// INFO: info: to(C[0:64])[256]
+// INFO: info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size={{[0-9]+}}, Name=DeviceEnvironment
// INFO: info: Creating new map entry with HstPtrBase={{.*}}, HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, DynRefCount=1, HoldRefCount=0, Name=A[0:64]
// INFO: info: Creating new map entry with HstPtrBase={{.*}}, HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, DynRefCount=0, HoldRefCount=1, Name=B[0:64]
-// INFO: info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=256, Name=B[0:64]
+// INFO: info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=256, Name=B[0:64], at info.c:{{[0-9]+}}:{{[0-9]+}}
// INFO: info: Creating new map entry with HstPtrBase={{.*}}, HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, DynRefCount=1, HoldRefCount=0, Name=C[0:64]
-// INFO: info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=256, Name=C[0:64]
+// INFO: info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=256, Name=C[0:64], at {{.*}}info.c:{{[0-9]+}}:{{[0-9]+}}
+// INFO: info: Waiting for asynchronous operations to complete at info.c:{{[0-9]+}}:{{[0-9]+}}
// INFO: info: OpenMP Host-Device pointer mappings after block at {{.*}}info.c:{{[0-9]+}}:{{[0-9]+}}:
// INFO: info: Host Ptr Target Ptr Size (B) DynRefCount HoldRefCount Declaration
// INFO: info: {{.*}} {{.*}} 256 1 0 C[0:64] at {{.*}}info.c:{{[0-9]+}}:{{[0-9]+}}
@@ -44,6 +46,7 @@ int main() {
// INFO: info: firstprivate(val)[4]
// INFO: info: Launching kernel __omp_offloading_{{.*}}main{{.*}} with [{{[0-9]+}},1,1] blocks and [{{[0-9]+}},1,1] threads in Generic mode
// AMDGPU: AMDGPU device {{[0-9]}} info: #Args: {{[0-9]}} Teams x Thrds: {{[0-9]+}}x {{[0-9]+}} (MaxFlatWorkGroupSize: {{[0-9]+}}) LDS Usage: {{[0-9]+}}B #SGPRs/VGPRs: {{[0-9]+}}/{{[0-9]+}} #SGPR/VGPR Spills: {{[0-9]+}}/{{[0-9]+}} Tripcount: {{[0-9]+}}
+// INFO: info: Waiting for asynchronous operations to complete at info.c:{{[0-9]+}}:{{[0-9]+}}
// INFO: info: OpenMP Host-Device pointer mappings after block at {{.*}}info.c:{{[0-9]+}}:{{[0-9]+}}:
// INFO: info: Host Ptr Target Ptr Size (B) DynRefCount HoldRefCount Declaration
// INFO: info: {{.*}} {{.*}} 256 1 0 C[0:64] at {{.*}}info.c:{{[0-9]+}}:{{[0-9]+}}
@@ -53,7 +56,8 @@ int main() {
// INFO: info: alloc(A[0:64])[256]
// INFO: info: tofrom(B[0:64])[256]
// INFO: info: to(C[0:64])[256]
-// INFO: info: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}}, Size=256, Name=B[0:64]
+// INFO: info: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}}, Size=256, Name=B[0:64], at info.c:{{[0-9]+}}:{{[0-9]+}}
+// INFO: info: Waiting for asynchronous operations to complete at info.c:{{[0-9]+}}:{{[0-9]+}}
// INFO: info: Removing map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, Name=C[0:64]
// INFO: info: Removing map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, Name=B[0:64]
// INFO: info: Removing map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, Name=A[0:64]
>From 2b1b7e4f5bd94fbcaaff2de51b029236a7c54818 Mon Sep 17 00:00:00 2001
From: Jason Van Beusekom <jason.van-beusekom at hpe.com>
Date: Wed, 16 Sep 2026 16:45:43 -0500
Subject: [PATCH 2/3] feedback
---
offload/include/Shared/SourceInfo.h | 13 -------
offload/libompaccsupport/Mapping.cpp | 8 +++--
offload/libomptarget/interface.cpp | 4 +--
offload/libomptarget/private.h | 10 ++++--
offload/test/offloading/info_nowait.c | 50 +++++++++++++++++++++++++++
5 files changed, 65 insertions(+), 20 deletions(-)
create mode 100644 offload/test/offloading/info_nowait.c
diff --git a/offload/include/Shared/SourceInfo.h b/offload/include/Shared/SourceInfo.h
index 2b42dec37549f..e8dfb46c973e0 100644
--- a/offload/include/Shared/SourceInfo.h
+++ b/offload/include/Shared/SourceInfo.h
@@ -99,17 +99,4 @@ static inline std::string getNameFromMapping(const map_var_info_t Name) {
return NameStr.substr(Begin + 1, End - Begin - 1);
}
-/// Returns "<Prefix><filename>:<line>:<column>" when \p Loc carries a
-/// compiler-provided source location, or an empty string otherwise, so info
-/// output can be annotated without emitting "unknown:0:0".
-static inline std::string getSourceLocationSuffix(const ident_t *Loc,
- const char *Prefix) {
- SourceInfo Info(Loc);
- if (!Info.isAvailible())
- return "";
- return std::string(Prefix) + Info.getFilename() + ":" +
- std::to_string(Info.getLine()) + ":" +
- std::to_string(Info.getColumn());
-}
-
#endif // OMPTARGET_SHARED_SOURCE_INFO_H
diff --git a/offload/libompaccsupport/Mapping.cpp b/offload/libompaccsupport/Mapping.cpp
index a4905e4106deb..d02689783c6c3 100644
--- a/offload/libompaccsupport/Mapping.cpp
+++ b/offload/libompaccsupport/Mapping.cpp
@@ -560,8 +560,12 @@ static void printCopyInfoImpl(int DeviceId, bool H2D, void *SrcPtrBegin,
void *DstPtrBegin, int64_t Size,
HostDataToTargetTy *HT, const ident_t *Loc,
const char *Name) {
-
- std::string LocStr = getSourceLocationSuffix(Loc, ", at ");
+ SourceInfo Info(Loc);
+ std::string LocStr;
+ if (Info.isAvailible())
+ LocStr = ", at " + std::string(Info.getFilename()) + ":" +
+ std::to_string(Info.getLine()) + ":" +
+ std::to_string(Info.getColumn());
INFO(OMP_INFOTYPE_DATA_TRANSFER, DeviceId,
"Copying data from %s to %s, %sPtr=" DPxMOD ", %sPtr=" DPxMOD
diff --git a/offload/libomptarget/interface.cpp b/offload/libomptarget/interface.cpp
index 51e7b6109c094..06a6ba56ffa61 100644
--- a/offload/libomptarget/interface.cpp
+++ b/offload/libomptarget/interface.cpp
@@ -187,7 +187,7 @@ targetData(ident_t *Loc, int64_t DeviceId, int32_t ArgNum, void **ArgsBase,
Rc = processAttachEntries(*DeviceOrErr, *StateInfo, AsyncInfo);
if (Rc == OFFLOAD_SUCCESS) {
- printSyncInfo(Loc, DeviceId);
+ printSyncInfo(Loc, DeviceId, AsyncInfo);
Rc = AsyncInfo.synchronize();
}
}
@@ -442,7 +442,7 @@ static inline int targetKernel(ident_t *Loc, int64_t DeviceId, int32_t NumTeams,
{ // required to show synchronization
TIMESCOPE_WITH_DETAILS_AND_IDENT("Runtime: synchronize", "", Loc);
if (Rc == OFFLOAD_SUCCESS) {
- printSyncInfo(Loc, DeviceId);
+ printSyncInfo(Loc, DeviceId, AsyncInfo);
Rc = AsyncInfo.synchronize();
}
diff --git a/offload/libomptarget/private.h b/offload/libomptarget/private.h
index 0de02cb2304d6..79bdffab9a2e0 100644
--- a/offload/libomptarget/private.h
+++ b/offload/libomptarget/private.h
@@ -89,12 +89,16 @@ printKernelArguments(const ident_t *Loc, const int64_t DeviceId,
////////////////////////////////////////////////////////////////////////////////
/// Report that the runtime is about to wait for the region's outstanding
/// asynchronous operations (data transfers and kernels) to complete.
-static inline void printSyncInfo(const ident_t *Loc, const int64_t DeviceId) {
+static inline void printSyncInfo(const ident_t *Loc, const int64_t DeviceId,
+ const AsyncInfoTy &AsyncInfo) {
if (!(getInfoLevel() & OMP_INFOTYPE_DATA_TRANSFER))
return;
- std::string LocStr = getSourceLocationSuffix(Loc, " at ");
+ if (AsyncInfo.SyncType != AsyncInfoTy::SyncTy::BLOCKING)
+ return;
+ SourceInfo Info(Loc);
INFO(OMP_INFOTYPE_DATA_TRANSFER, DeviceId,
- "Waiting for asynchronous operations to complete%s\n", LocStr.c_str());
+ "Waiting for asynchronous operations to complete at %s:%d:%d\n",
+ Info.getFilename(), Info.getLine(), Info.getColumn());
}
////////////////////////////////////////////////////////////////////////////////
diff --git a/offload/test/offloading/info_nowait.c b/offload/test/offloading/info_nowait.c
new file mode 100644
index 0000000000000..af2cbefe0151b
--- /dev/null
+++ b/offload/test/offloading/info_nowait.c
@@ -0,0 +1,50 @@
+// Verify LIBOMPTARGET_INFO reporting of data-transfer synchronization.
+//
+// RUN: %libomptarget-compile-generic -gline-tables-only -fopenmp-extensions
+// RUN: env LIBOMPTARGET_INFO=32 %libomptarget-run-generic 2>&1 | \
+// RUN: %fcheck-generic -allow-empty -check-prefixes=INFO
+
+// FIXME: Fails due to optimized debugging in 'ptxas'.
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+
+#include <stdio.h>
+
+int main(void) {
+ int x = 0, y = 0;
+
+ // Blocking region: synchronization is blocking, so the wait is announced.
+#pragma omp target map(tofrom : x)
+ x = 1;
+
+ // Nowait kernel inside an active parallel/single: it has a task team and
+ // therefore synchronizes non-blockingly, so no wait must be announced even
+ // though its data transfers are still reported.
+#pragma omp parallel num_threads(2)
+#pragma omp single
+ {
+#pragma omp target map(tofrom : x) nowait
+ x = 2;
+#pragma omp taskwait
+ }
+
+ // Nowait target-data constructs reach the synchronization notice through a
+ // different runtime entry point; they are non-blocking here too and likewise
+ // must not announce a wait.
+#pragma omp parallel num_threads(2)
+#pragma omp single
+ {
+#pragma omp target enter data map(to : y) nowait
+#pragma omp taskwait
+#pragma omp target exit data map(from : y) nowait
+#pragma omp taskwait
+ }
+
+ printf("x = %d, y = %d\n", x, y);
+ return x != 2;
+}
+
+// clang-format off
+// INFO: info: Waiting for asynchronous operations to complete at info_nowait.c:{{[0-9]+}}:{{[0-9]+}}
+// INFO: info: Copying data from host to device,{{.*}}at info_nowait.c:{{[0-9]+}}:{{[0-9]+}}
+// INFO-NOT: Waiting for asynchronous operations to complete
+// clang-format on
>From 0a72fef9ed46fe4dc6fee05e6075fe8dc8347da3 Mon Sep 17 00:00:00 2001
From: Jason Van Beusekom <jason.van-beusekom at hpe.com>
Date: Tue, 29 Sep 2026 14:57:17 -0500
Subject: [PATCH 3/3] fix broken tests
---
offload/test/offloading/info.c | 12 ++++++------
offload/test/offloading/info_nowait.c | 4 ++--
2 files changed, 8 insertions(+), 8 deletions(-)
diff --git a/offload/test/offloading/info.c b/offload/test/offloading/info.c
index cfe4ad3afaa10..7807193624dc3 100644
--- a/offload/test/offloading/info.c
+++ b/offload/test/offloading/info.c
@@ -30,13 +30,13 @@ int main() {
// INFO: info: alloc(A[0:64])[256]
// INFO: info: tofrom(B[0:64])[256]
// INFO: info: to(C[0:64])[256]
-// INFO: info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size={{[0-9]+}}, Name=DeviceEnvironment
+// AMDGPU: info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size={{[0-9]+}}, Name=DeviceEnvironment
// INFO: info: Creating new map entry with HstPtrBase={{.*}}, HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, DynRefCount=1, HoldRefCount=0, Name=A[0:64]
// INFO: info: Creating new map entry with HstPtrBase={{.*}}, HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, DynRefCount=0, HoldRefCount=1, Name=B[0:64]
-// INFO: info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=256, Name=B[0:64], at info.c:{{[0-9]+}}:{{[0-9]+}}
+// INFO: info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=256, Name=B[0:64], at {{.*}}info.c:{{[0-9]+}}:{{[0-9]+}}
// INFO: info: Creating new map entry with HstPtrBase={{.*}}, HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, DynRefCount=1, HoldRefCount=0, Name=C[0:64]
// INFO: info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=256, Name=C[0:64], at {{.*}}info.c:{{[0-9]+}}:{{[0-9]+}}
-// INFO: info: Waiting for asynchronous operations to complete at info.c:{{[0-9]+}}:{{[0-9]+}}
+// INFO: info: Waiting for asynchronous operations to complete at {{.*}}info.c:{{[0-9]+}}:{{[0-9]+}}
// INFO: info: OpenMP Host-Device pointer mappings after block at {{.*}}info.c:{{[0-9]+}}:{{[0-9]+}}:
// INFO: info: Host Ptr Target Ptr Size (B) DynRefCount HoldRefCount Declaration
// INFO: info: {{.*}} {{.*}} 256 1 0 C[0:64] at {{.*}}info.c:{{[0-9]+}}:{{[0-9]+}}
@@ -46,7 +46,7 @@ int main() {
// INFO: info: firstprivate(val)[4]
// INFO: info: Launching kernel __omp_offloading_{{.*}}main{{.*}} with [{{[0-9]+}},1,1] blocks and [{{[0-9]+}},1,1] threads in Generic mode
// AMDGPU: AMDGPU device {{[0-9]}} info: #Args: {{[0-9]}} Teams x Thrds: {{[0-9]+}}x {{[0-9]+}} (MaxFlatWorkGroupSize: {{[0-9]+}}) LDS Usage: {{[0-9]+}}B #SGPRs/VGPRs: {{[0-9]+}}/{{[0-9]+}} #SGPR/VGPR Spills: {{[0-9]+}}/{{[0-9]+}} Tripcount: {{[0-9]+}}
-// INFO: info: Waiting for asynchronous operations to complete at info.c:{{[0-9]+}}:{{[0-9]+}}
+// INFO: info: Waiting for asynchronous operations to complete at {{.*}}info.c:{{[0-9]+}}:{{[0-9]+}}
// INFO: info: OpenMP Host-Device pointer mappings after block at {{.*}}info.c:{{[0-9]+}}:{{[0-9]+}}:
// INFO: info: Host Ptr Target Ptr Size (B) DynRefCount HoldRefCount Declaration
// INFO: info: {{.*}} {{.*}} 256 1 0 C[0:64] at {{.*}}info.c:{{[0-9]+}}:{{[0-9]+}}
@@ -56,8 +56,8 @@ int main() {
// INFO: info: alloc(A[0:64])[256]
// INFO: info: tofrom(B[0:64])[256]
// INFO: info: to(C[0:64])[256]
-// INFO: info: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}}, Size=256, Name=B[0:64], at info.c:{{[0-9]+}}:{{[0-9]+}}
-// INFO: info: Waiting for asynchronous operations to complete at info.c:{{[0-9]+}}:{{[0-9]+}}
+// INFO: info: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}}, Size=256, Name=B[0:64], at {{.*}}info.c:{{[0-9]+}}:{{[0-9]+}}
+// INFO: info: Waiting for asynchronous operations to complete at {{.*}}info.c:{{[0-9]+}}:{{[0-9]+}}
// INFO: info: Removing map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, Name=C[0:64]
// INFO: info: Removing map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, Name=B[0:64]
// INFO: info: Removing map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, Name=A[0:64]
diff --git a/offload/test/offloading/info_nowait.c b/offload/test/offloading/info_nowait.c
index af2cbefe0151b..bec706687629a 100644
--- a/offload/test/offloading/info_nowait.c
+++ b/offload/test/offloading/info_nowait.c
@@ -44,7 +44,7 @@ int main(void) {
}
// clang-format off
-// INFO: info: Waiting for asynchronous operations to complete at info_nowait.c:{{[0-9]+}}:{{[0-9]+}}
-// INFO: info: Copying data from host to device,{{.*}}at info_nowait.c:{{[0-9]+}}:{{[0-9]+}}
+// INFO: info: Waiting for asynchronous operations to complete at {{.*}}info_nowait.c:{{[0-9]+}}:{{[0-9]+}}
+// INFO: info: Copying data from host to device,{{.*}}at {{.*}}info_nowait.c:{{[0-9]+}}:{{[0-9]+}}
// INFO-NOT: Waiting for asynchronous operations to complete
// clang-format on
More information about the llvm-commits
mailing list