[llvm] [OpenMP][Offload] Handle `present/to/from` when a different entry did `alloc/delete`. (PR #165494)
Abhinav Gaba via llvm-commits
llvm-commits at lists.llvm.org
Tue Feb 10 06:43:26 PST 2026
https://github.com/abhinavgaba updated https://github.com/llvm/llvm-project/pull/165494
>From 8f583973a9692e00122cee93cb79a3cc730a8f6a Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Tue, 28 Oct 2025 14:16:23 -0700
Subject: [PATCH 01/16] [OpenMP][Offload] Handle for non-memberof
present/to/from entries irrespective of order.
For cases like:
```c
map(alloc: x) map(to: x)
```
If the entry of `map(to: x)` is encountered after the entry for
`map(alloc:x)`, we still want to do a data-transfer even though the
ref-count of `x` was already 0, because the new allocation for `x`
happened as part of the current directive.
Similarly, for:
```c
... map(alloc: x) map(from: x)
```
If the entry for `map(from:x)` is encountered before the entry for
`map(alloc:x)`, we want to do a data-transfer even though the
ref-count was not 0 when looking at the `from` entry, because by the end of
the directive, the ref-count of `x` will go down to zero.
And for:
```c
... map(from : x) map(alloc, present: x)
```
If the "present" entry is encountered after the "from" entry, then it becomes
a no-op, as the "from" entry will do an allocation if no match was found.
In this PR, these are handled by the runtime via the following:
* For `to` and `present`, we also look-up in the existing table where we tracked
new allocations when making the decision for the entry.
* For `from`, we keep track of any deferred data transfers and when the
ref-count of a pointer goes to zero, see if there were any previously
deferred `from` transfers for that pointer.
This can be done in the compiler, and that would avoid any runtime
overhead, but it would require creating two separate offload struct entries
for the entry and exit mappings (even for the `target` construct),
with properly decayed maps, and either:
(1) sorted in order of:
* `present > to > ...` for the implied `target enter data`; and
* `from > ...` for the `target exit data`
e.g.
```c
#pragma omp target map(to: x) map(present, alloc: x) map(always, from: x)
// has to be broken into:
// from becomes alloc on entry:
// #pragma omp target enter data map(present, alloc: x)
// map(to: x)
// map(alloc: x)
//
// "present" and "to" just "decay" into "alloc"
// #pragma omp target exit data map(always, from: x)
// map(alloc: x)
// map(alloc: x)
```
Or,
(2) Merged into one entry each on the `target enter/exit data`
directives.
```c
#pragma omp target map(to: x) map(present, alloc: x) map(always, from: x)
// has to be broken into:
// from becomes alloc on entry:
// #pragma omp target enter data map(present, to: x)
//
// "present" and "to" just "decay" into "alloc"
// #pragma omp target exit data map(always, from: x)
```
The number of entries on the two would need to stay the same on the two to avoid
ref-count mismatch.
(1) would be simpler, but won't likely work for cases like:
```c
... map(delete: x) map(from:x)
```
as there is no clear "winner" between the two. So, for such cases, the compiler
would likely have to do (2), which is the cleanest solution, but will take
longer to implement. For EXPR comparisons, it can build-upon the
`AttachPtrExprComparator` that was implemented as part of #153683,
but that should probably wait for the PR to be merged to avoid
conflicts.
Another alternative is to sort the entries in the runtime, which may be
slower than on-demand lookups/updates that this PR does, because we
always would be doing this sorting even when not needed, but may be
faster in others where the constant-time overhead of map/set
insertions/lookups becomes too large because of the number of maps. But
that will still have to worry about the `from` + `delete` case.
---
offload/include/OpenMP/Mapping.h | 34 +++--
offload/libomptarget/OpenMP/Mapping.cpp | 15 +-
offload/libomptarget/interface.cpp | 17 ++-
offload/libomptarget/omptarget.cpp | 137 +++++++++++++-----
.../mapping/map_ordering_tgt_alloc_from_to.c | 14 ++
.../map_ordering_tgt_alloc_present_tofrom.c | 27 ++++
.../mapping/map_ordering_tgt_alloc_tofrom.c | 14 ++
.../map_ordering_tgt_data_alloc_from.c | 14 ++
.../map_ordering_tgt_data_alloc_to_from.c | 17 +++
.../map_ordering_tgt_data_alloc_tofrom.c | 17 +++
10 files changed, 249 insertions(+), 57 deletions(-)
create mode 100644 offload/test/mapping/map_ordering_tgt_alloc_from_to.c
create mode 100644 offload/test/mapping/map_ordering_tgt_alloc_present_tofrom.c
create mode 100644 offload/test/mapping/map_ordering_tgt_alloc_tofrom.c
create mode 100644 offload/test/mapping/map_ordering_tgt_data_alloc_from.c
create mode 100644 offload/test/mapping/map_ordering_tgt_data_alloc_to_from.c
create mode 100644 offload/test/mapping/map_ordering_tgt_data_alloc_tofrom.c
diff --git a/offload/include/OpenMP/Mapping.h b/offload/include/OpenMP/Mapping.h
index 45bd9c6e7da8b..517f6c0a99244 100644
--- a/offload/include/OpenMP/Mapping.h
+++ b/offload/include/OpenMP/Mapping.h
@@ -484,20 +484,26 @@ struct AttachMapInfo {
MapType(Type), Pointername(Name) {}
};
-/// Structure to track ATTACH entries and new allocations across recursive calls
-/// (for handling mappers) to targetDataBegin for a given construct.
-struct AttachInfoTy {
- /// ATTACH map entries for deferred processing.
+/// Structure to track new allocations, ATTACH entries and deferred data
+/// transfer information for a given construct, across recursive calls (for
+/// handling mappers) to targetDataBegin/targetDataEnd.
+struct StateInfoTy {
+ /// ATTACH map entries for deferred processing until all other maps are done.
llvm::SmallVector<AttachMapInfo> AttachEntries;
+ /// Host pointers for which new memory was allocated.
/// Key: host pointer, Value: allocation size.
llvm::DenseMap<void *, int64_t> NewAllocations;
- AttachInfoTy() = default;
+ /// Host pointers that had a FROM entry, but for which a data transfer didn't
+ /// occur due to the ref-count not being zero.
+ llvm::SmallSet<void *, 32> DeferredFromPtrs;
+
+ StateInfoTy() = default;
// Delete copy constructor and copy assignment operator to prevent copying
- AttachInfoTy(const AttachInfoTy &) = delete;
- AttachInfoTy &operator=(const AttachInfoTy &) = delete;
+ StateInfoTy(const StateInfoTy &) = delete;
+ StateInfoTy &operator=(const StateInfoTy &) = delete;
};
// Function pointer type for targetData* functions (targetDataBegin,
@@ -505,7 +511,7 @@ struct AttachInfoTy {
typedef int (*TargetDataFuncPtrTy)(ident_t *, DeviceTy &, int32_t, void **,
void **, int64_t *, int64_t *,
map_var_info_t *, void **, AsyncInfoTy &,
- AttachInfoTy *, bool);
+ StateInfoTy *, bool);
void dumpTargetPointerMappings(const ident_t *Loc, DeviceTy &Device,
bool toStdOut = false);
@@ -514,24 +520,22 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
void **ArgsBase, void **Args, int64_t *ArgSizes,
int64_t *ArgTypes, map_var_info_t *ArgNames,
void **ArgMappers, AsyncInfoTy &AsyncInfo,
- AttachInfoTy *AttachInfo = nullptr,
- bool FromMapper = false);
+ StateInfoTy *StateInfo = nullptr, bool FromMapper = false);
int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
void **ArgBases, void **Args, int64_t *ArgSizes,
int64_t *ArgTypes, map_var_info_t *ArgNames,
void **ArgMappers, AsyncInfoTy &AsyncInfo,
- AttachInfoTy *AttachInfo = nullptr, bool FromMapper = false);
+ StateInfoTy *StateInfo = nullptr, bool FromMapper = false);
int targetDataUpdate(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
void **ArgsBase, void **Args, int64_t *ArgSizes,
int64_t *ArgTypes, map_var_info_t *ArgNames,
void **ArgMappers, AsyncInfoTy &AsyncInfo,
- AttachInfoTy *AttachInfo = nullptr,
- bool FromMapper = false);
+ StateInfoTy *StateInfo = nullptr, bool FromMapper = false);
// Process deferred ATTACH map entries collected during targetDataBegin.
-int processAttachEntries(DeviceTy &Device, AttachInfoTy &AttachInfo,
+int processAttachEntries(DeviceTy &Device, StateInfoTy &StateInfo,
AsyncInfoTy &AsyncInfo);
struct MappingInfoTy {
@@ -572,7 +576,7 @@ struct MappingInfoTy {
bool HasFlagTo, bool HasFlagAlways, bool IsImplicit, bool UpdateRefCount,
bool HasCloseModifier, bool HasPresentModifier, bool HasHoldModifier,
AsyncInfoTy &AsyncInfo, HostDataToTargetTy *OwnedTPR = nullptr,
- bool ReleaseHDTTMap = true);
+ bool ReleaseHDTTMap = true, StateInfoTy *StateInfo = nullptr);
/// Return the target pointer for \p HstPtrBegin in \p HDTTMap. The accessor
/// ensures exclusive access to the HDTT map.
diff --git a/offload/libomptarget/OpenMP/Mapping.cpp b/offload/libomptarget/OpenMP/Mapping.cpp
index 9b3533895f2a6..a3f634bc0a9eb 100644
--- a/offload/libomptarget/OpenMP/Mapping.cpp
+++ b/offload/libomptarget/OpenMP/Mapping.cpp
@@ -202,7 +202,8 @@ TargetPointerResultTy MappingInfoTy::getTargetPointer(
int64_t TgtPadding, int64_t Size, map_var_info_t HstPtrName, bool HasFlagTo,
bool HasFlagAlways, bool IsImplicit, bool UpdateRefCount,
bool HasCloseModifier, bool HasPresentModifier, bool HasHoldModifier,
- AsyncInfoTy &AsyncInfo, HostDataToTargetTy *OwnedTPR, bool ReleaseHDTTMap) {
+ AsyncInfoTy &AsyncInfo, HostDataToTargetTy *OwnedTPR, bool ReleaseHDTTMap,
+ StateInfoTy *StateInfo) {
LookupResult LR = lookupMapping(HDTTMap, HstPtrBegin, Size, OwnedTPR);
LR.TPR.Flags.IsPresent = true;
@@ -324,8 +325,18 @@ TargetPointerResultTy MappingInfoTy::getTargetPointer(
// If the target pointer is valid, and we need to transfer data, issue the
// data transfer.
+ auto WasNewlyAllocatedOnCurrentConstruct = [&]() {
+ if (!StateInfo)
+ return false;
+ return StateInfo->NewAllocations.contains(HstPtrBegin);
+ };
+
+ // Even if this isn't a new entry, we still need to do a data-transfer if
+ // the pointer was newly allocated previously on the same construct.
if (LR.TPR.TargetPointer && !LR.TPR.Flags.IsHostPointer && HasFlagTo &&
- (LR.TPR.Flags.IsNewEntry || HasFlagAlways) && Size != 0) {
+ (LR.TPR.Flags.IsNewEntry || HasFlagAlways ||
+ WasNewlyAllocatedOnCurrentConstruct()) &&
+ Size != 0) {
// If we have something like:
// #pragma omp target map(to: s.myarr[0:10]) map(to: s.myarr[0:10])
diff --git a/offload/libomptarget/interface.cpp b/offload/libomptarget/interface.cpp
index fe18289765906..ac03546860740 100644
--- a/offload/libomptarget/interface.cpp
+++ b/offload/libomptarget/interface.cpp
@@ -167,19 +167,22 @@ targetData(ident_t *Loc, int64_t DeviceId, int32_t ArgNum, void **ArgsBase,
int Rc = OFFLOAD_SUCCESS;
- // Only allocate AttachInfo for targetDataBegin
- std::unique_ptr<AttachInfoTy> AttachInfo;
- if (TargetDataFunction == targetDataBegin)
- AttachInfo = std::make_unique<AttachInfoTy>();
+ // Allocate StateInfo for targetDataBegin and targetDataEnd to track
+ // allocations, pointer attachments and deferred transfers.
+ // This is not needed for targetDataUpdate.
+ std::unique_ptr<StateInfoTy> StateInfo;
+ if (TargetDataFunction == targetDataBegin ||
+ TargetDataFunction == targetDataEnd)
+ StateInfo = std::make_unique<StateInfoTy>();
Rc = TargetDataFunction(Loc, *DeviceOrErr, ArgNum, ArgsBase, Args, ArgSizes,
ArgTypes, ArgNames, ArgMappers, AsyncInfo,
- AttachInfo.get(), /*FromMapper=*/false);
+ StateInfo.get(), /*FromMapper=*/false);
if (Rc == OFFLOAD_SUCCESS) {
// Process deferred ATTACH entries BEFORE synchronization
- if (AttachInfo && !AttachInfo->AttachEntries.empty())
- Rc = processAttachEntries(*DeviceOrErr, *AttachInfo, AsyncInfo);
+ if (StateInfo && !StateInfo->AttachEntries.empty())
+ Rc = processAttachEntries(*DeviceOrErr, *StateInfo, AsyncInfo);
if (Rc == OFFLOAD_SUCCESS)
Rc = AsyncInfo.synchronize();
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index 69725e77bae00..bef1488b2956f 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -294,7 +294,7 @@ int targetDataMapper(ident_t *Loc, DeviceTy &Device, void *ArgBase, void *Arg,
int64_t ArgSize, int64_t ArgType, map_var_info_t ArgNames,
void *ArgMapper, AsyncInfoTy &AsyncInfo,
TargetDataFuncPtrTy TargetDataFunction,
- AttachInfoTy *AttachInfo = nullptr) {
+ StateInfoTy *StateInfo = nullptr) {
DP("Calling the mapper function " DPxMOD "\n", DPxPTR(ArgMapper));
// The mapper function fills up Components.
@@ -325,7 +325,7 @@ int targetDataMapper(ident_t *Loc, DeviceTy &Device, void *ArgBase, void *Arg,
MapperArgsBase.data(), MapperArgs.data(),
MapperArgSizes.data(), MapperArgTypes.data(),
MapperArgNames.data(), /*arg_mappers*/ nullptr,
- AsyncInfo, AttachInfo, /*FromMapper=*/true);
+ AsyncInfo, StateInfo, /*FromMapper=*/true);
return Rc;
}
@@ -509,12 +509,12 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
void **ArgsBase, void **Args, int64_t *ArgSizes,
int64_t *ArgTypes, map_var_info_t *ArgNames,
void **ArgMappers, AsyncInfoTy &AsyncInfo,
- AttachInfoTy *AttachInfo, bool FromMapper) {
- assert(AttachInfo && "AttachInfo must be available for targetDataBegin for "
- "handling ATTACH map-types.");
+ StateInfoTy *StateInfo, bool FromMapper) {
+ assert(StateInfo && "StateInfo must be available for targetDataBegin for "
+ "handling ATTACH and TO/TOFROM map-types.");
// process each input.
for (int32_t I = 0; I < ArgNum; ++I) {
- // Ignore private variables and arrays - there is no mapping for them.
+ // Ignore private variables and arrays - there is no mapping for t.attahem.
if ((ArgTypes[I] & OMP_TGT_MAPTYPE_LITERAL) ||
(ArgTypes[I] & OMP_TGT_MAPTYPE_PRIVATE))
continue;
@@ -529,7 +529,7 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
map_var_info_t ArgName = (!ArgNames) ? nullptr : ArgNames[I];
int Rc = targetDataMapper(Loc, Device, ArgsBase[I], Args[I], ArgSizes[I],
ArgTypes[I], ArgName, ArgMappers[I], AsyncInfo,
- targetDataBegin, AttachInfo);
+ targetDataBegin, StateInfo);
if (Rc != OFFLOAD_SUCCESS) {
REPORT("Call to targetDataBegin via targetDataMapper for custom mapper"
@@ -556,7 +556,7 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
// similar to firstprivate (PRIVATE | TO) entries by
// PrivateArgumentManager.
if (!IsCorrespondingPointerInit)
- AttachInfo->AttachEntries.emplace_back(
+ StateInfo->AttachEntries.emplace_back(
/*PointerBase=*/HstPtrBase, /*PointeeBegin=*/HstPtrBegin,
/*PointerSize=*/DataSize, /*MapType=*/ArgTypes[I],
/*PointeeName=*/HstPtrName);
@@ -633,7 +633,7 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
// Track new allocation, for eventual use in attachment decision-making.
if (PointerTpr.Flags.IsNewEntry && !IsHostPtr)
- AttachInfo->NewAllocations[HstPtrBase] = sizeof(void *);
+ StateInfo->NewAllocations[HstPtrBase] = sizeof(void *);
DP("There are %zu bytes allocated at target address " DPxMOD " - is%s new"
"\n",
@@ -654,7 +654,8 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
auto TPR = Device.getMappingInfo().getTargetPointer(
HDTTMap, HstPtrBegin, HstPtrBase, TgtPadding, DataSize, HstPtrName,
HasFlagTo, HasFlagAlways, IsImplicit, UpdateRef, HasCloseModifier,
- HasPresentModifier, HasHoldModifier, AsyncInfo, PointerTpr.getEntry());
+ HasPresentModifier, HasHoldModifier, AsyncInfo, PointerTpr.getEntry(),
+ /*ReleaseHDTTMap=*/true, StateInfo);
void *TgtPtrBegin = TPR.TargetPointer;
IsHostPtr = TPR.Flags.IsHostPointer;
// If data_size==0, then the argument could be a zero-length pointer to
@@ -664,11 +665,30 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
HasPresentModifier ? "'present' map type modifier"
: "device failure or illegal mapping");
return OFFLOAD_FAIL;
+ } else if (TgtPtrBegin && HasPresentModifier &&
+ StateInfo->NewAllocations.contains(HstPtrBegin)) {
+ // For "PRESENT" entries, we may have cases like the following:
+ // map(alloc: p[0]) map(present, alloc: p[0])
+ // If the compiler does not merge these entries, then the "PRESENT" entry
+ // may be encountered after a previous entry allocated new storage for it.
+ // To catch such cases, we should also look at any existing allocations
+ // and error out if we have one matching the pointer. We don't need to
+ // worry about cases like:
+ // map(alloc: p[1:10]) map(present, alloc: p[2:5])
+ // as the list-items share storage, but are not identical, which is a
+ // user error as per OpenMP.
+ MESSAGE("device mapping required by 'present' map type modifier does not "
+ "exist for host address " DPxMOD " (%" PRId64 " bytes)\n",
+ DPxPTR(HstPtrBegin), DataSize);
+ REPORT("Pointer " DPxMOD
+ " was not present on the device upon entry to the region.\n",
+ DPxPTR(HstPtrBegin));
+ return OFFLOAD_FAIL;
}
- // Track new allocation, for eventual use in attachment decision-making.
+ // Track new allocation, for eventual use in attachment/to decision-making.
if (TPR.Flags.IsNewEntry && !IsHostPtr && TgtPtrBegin)
- AttachInfo->NewAllocations[HstPtrBegin] = DataSize;
+ StateInfo->NewAllocations[HstPtrBegin] = DataSize;
DP("There are %" PRId64 " bytes allocated at target address " DPxMOD
" - is%s new\n",
@@ -751,29 +771,29 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
///
/// For this purpose, we insert a data_fence before the first
/// pointer-attachment, (3), to ensure that all pending transfers finish first.
-int processAttachEntries(DeviceTy &Device, AttachInfoTy &AttachInfo,
+int processAttachEntries(DeviceTy &Device, StateInfoTy &StateInfo,
AsyncInfoTy &AsyncInfo) {
// Report all tracked allocations from both main loop and ATTACH processing
- if (!AttachInfo.NewAllocations.empty()) {
+ if (!StateInfo.NewAllocations.empty()) {
DP("Tracked %u total new allocations:\n",
- (unsigned)AttachInfo.NewAllocations.size());
- for ([[maybe_unused]] const auto &Alloc : AttachInfo.NewAllocations) {
+ (unsigned)StateInfo.NewAllocations.size());
+ for ([[maybe_unused]] const auto &Alloc : StateInfo.NewAllocations) {
DP(" Host ptr: " DPxMOD ", Size: %" PRId64 " bytes\n",
DPxPTR(Alloc.first), Alloc.second);
}
}
- if (AttachInfo.AttachEntries.empty())
+ if (StateInfo.AttachEntries.empty())
return OFFLOAD_SUCCESS;
DP("Processing %zu deferred ATTACH map entries\n",
- AttachInfo.AttachEntries.size());
+ StateInfo.AttachEntries.size());
int Ret = OFFLOAD_SUCCESS;
bool IsFirstPointerAttachment = true;
- for (size_t EntryIdx = 0; EntryIdx < AttachInfo.AttachEntries.size();
+ for (size_t EntryIdx = 0; EntryIdx < StateInfo.AttachEntries.size();
++EntryIdx) {
- const auto &AttachEntry = AttachInfo.AttachEntries[EntryIdx];
+ const auto &AttachEntry = StateInfo.AttachEntries[EntryIdx];
void **HstPtr = reinterpret_cast<void **>(AttachEntry.PointerBase);
@@ -792,7 +812,7 @@ int processAttachEntries(DeviceTy &Device, AttachInfoTy &AttachInfo,
// Lambda to check if a pointer was newly allocated
auto WasNewlyAllocated = [&](void *Ptr, const char *PtrName) {
bool IsNewlyAllocated =
- llvm::any_of(AttachInfo.NewAllocations, [&](const auto &Alloc) {
+ llvm::any_of(StateInfo.NewAllocations, [&](const auto &Alloc) {
void *AllocPtr = Alloc.first;
int64_t AllocSize = Alloc.second;
return Ptr >= AllocPtr &&
@@ -1009,7 +1029,9 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
void **ArgBases, void **Args, int64_t *ArgSizes,
int64_t *ArgTypes, map_var_info_t *ArgNames,
void **ArgMappers, AsyncInfoTy &AsyncInfo,
- AttachInfoTy *AttachInfo, bool FromMapper) {
+ StateInfoTy *StateInfo, bool FromMapper) {
+ assert(StateInfo && "StateInfo is required for targetDataEnd for handling "
+ "FROM data transfers");
int Ret = OFFLOAD_SUCCESS;
auto *PostProcessingPtrs = new SmallVector<PostProcessingInfo>();
// process each input.
@@ -1037,7 +1059,7 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
map_var_info_t ArgName = (!ArgNames) ? nullptr : ArgNames[I];
Ret = targetDataMapper(Loc, Device, ArgBases[I], Args[I], ArgSizes[I],
ArgTypes[I], ArgName, ArgMappers[I], AsyncInfo,
- targetDataEnd);
+ targetDataEnd, StateInfo);
if (Ret != OFFLOAD_SUCCESS) {
REPORT("Call to targetDataEnd via targetDataMapper for custom mapper"
@@ -1106,8 +1128,28 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
// Move data back to the host
const bool HasAlways = ArgTypes[I] & OMP_TGT_MAPTYPE_ALWAYS;
const bool HasFrom = ArgTypes[I] & OMP_TGT_MAPTYPE_FROM;
- if (HasFrom && (HasAlways || TPR.Flags.IsLast) &&
- !TPR.Flags.IsHostPointer && DataSize != 0) {
+ const bool IsMemberOf = ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF;
+ // Lambda to check if there was a previously deferred FROM for this pointer
+ // due to its ref-count not being zero.
+ auto HasDeferredMapFrom = [&]() -> bool {
+ if (!StateInfo->DeferredFromPtrs.contains(HstPtrBegin))
+ return false;
+ DP("Found previously deferred FROM transfer for HstPtr=" DPxMOD "\n",
+ DPxPTR(HstPtrBegin));
+ // Remove it so we don't look at it again
+ StateInfo->DeferredFromPtrs.erase(HstPtrBegin);
+ return true;
+ };
+
+ bool IsMapFromOnNonHostNonZeroData =
+ HasFrom && !TPR.Flags.IsHostPointer && DataSize != 0;
+ bool IsLastOrHasAlways = TPR.Flags.IsLast || HasAlways;
+
+ if ((IsMapFromOnNonHostNonZeroData && IsLastOrHasAlways) ||
+ // Even if are not looking at an entry with FROM map-type, if there were
+ // any previously deferred FROM transfers for this pointer, we should
+ // do them when the ref-count goes down to zero.
+ (TPR.Flags.IsLast && HasDeferredMapFrom())) {
DP("Moving %" PRId64 " bytes (tgt:" DPxMOD ") -> (hst:" DPxMOD ")\n",
DataSize, DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin));
TIMESCOPE_WITH_DETAILS_AND_IDENT(
@@ -1137,6 +1179,30 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
OFFLOAD_SUCCESS)
return OFFLOAD_FAIL;
}
+ } else if (IsMapFromOnNonHostNonZeroData && !IsLastOrHasAlways &&
+ !IsMemberOf) {
+ // We can have cases like the following:
+ // map(alloc: p[0:1]) map(from: p[0:1])
+ //
+ // For such cases, if we have different entries for the two maps, we
+ // may not see the ref-count go down to zero when handling the From entry.
+ //
+ // So, we defer the FROM data-transfer until the ref-count goes down to
+ // zero (if it does).
+ //
+ // This should be limited to non-member-of entries because for member-of,
+ // their ref-count should go down only once as part of the parent.
+ //
+ // Also, we don't need to worry about cases like:
+ // map(alloc: p[0:10]) map(from: p[0:1])
+ //
+ // because that is not OpenMP 6.0 compliant, so we can just save the
+ // pointer without saving the size, and assume that the size for the
+ // "alloc" map will match that of "from".
+ StateInfo->DeferredFromPtrs.insert(HstPtrBegin);
+ DP("Deferring FROM map transfer for HstPtr=" DPxMOD ", Size=%" PRId64
+ "\n",
+ DPxPTR(HstPtrBegin), DataSize);
}
// Add pointer to the buffer for post-synchronize processing.
@@ -1315,7 +1381,7 @@ int targetDataUpdate(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
void **ArgsBase, void **Args, int64_t *ArgSizes,
int64_t *ArgTypes, map_var_info_t *ArgNames,
void **ArgMappers, AsyncInfoTy &AsyncInfo,
- AttachInfoTy *AttachInfo, bool FromMapper) {
+ StateInfoTy *StateInfo, bool FromMapper) {
// process each input.
for (int32_t I = 0; I < ArgNum; ++I) {
if ((ArgTypes[I] & OMP_TGT_MAPTYPE_LITERAL) ||
@@ -1806,21 +1872,21 @@ static int processDataBefore(ident_t *Loc, int64_t DeviceId, void *HostPtr,
if (!DeviceOrErr)
FATAL_MESSAGE(DeviceId, "%s", toString(DeviceOrErr.takeError()).c_str());
- // Create AttachInfo for tracking any ATTACH entries, or new-allocations
+ // Create StateInfo for tracking any ATTACH entries, new allocations,
// when handling the "begin" mapping for a target constructs.
- AttachInfoTy AttachInfo;
+ StateInfoTy StateInfo;
int Ret = targetDataBegin(Loc, *DeviceOrErr, ArgNum, ArgBases, Args, ArgSizes,
ArgTypes, ArgNames, ArgMappers, AsyncInfo,
- &AttachInfo, false /*FromMapper=*/);
+ &StateInfo, false /*FromMapper=*/);
if (Ret != OFFLOAD_SUCCESS) {
REPORT("Call to targetDataBegin failed, abort target.\n");
return OFFLOAD_FAIL;
}
// Process collected ATTACH entries
- if (!AttachInfo.AttachEntries.empty()) {
- Ret = processAttachEntries(*DeviceOrErr, AttachInfo, AsyncInfo);
+ if (!StateInfo.AttachEntries.empty()) {
+ Ret = processAttachEntries(*DeviceOrErr, StateInfo, AsyncInfo);
if (Ret != OFFLOAD_SUCCESS) {
REPORT("Failed to process ATTACH entries.\n");
return OFFLOAD_FAIL;
@@ -1987,9 +2053,14 @@ static int processDataAfter(ident_t *Loc, int64_t DeviceId, void *HostPtr,
if (!DeviceOrErr)
FATAL_MESSAGE(DeviceId, "%s", toString(DeviceOrErr.takeError()).c_str());
+ // Create StateInfo for tracking map(from)s for which ref-count is non-zero
+ // when the entry is encountered.
+ StateInfoTy StateInfo;
+
// Move data from device.
- int Ret = targetDataEnd(Loc, *DeviceOrErr, ArgNum, ArgBases, Args, ArgSizes,
- ArgTypes, ArgNames, ArgMappers, AsyncInfo);
+ int Ret =
+ targetDataEnd(Loc, *DeviceOrErr, ArgNum, ArgBases, Args, ArgSizes,
+ ArgTypes, ArgNames, ArgMappers, AsyncInfo, &StateInfo);
if (Ret != OFFLOAD_SUCCESS) {
REPORT("Call to targetDataEnd failed, abort target.\n");
return OFFLOAD_FAIL;
diff --git a/offload/test/mapping/map_ordering_tgt_alloc_from_to.c b/offload/test/mapping/map_ordering_tgt_alloc_from_to.c
new file mode 100644
index 0000000000000..67c88e7238842
--- /dev/null
+++ b/offload/test/mapping/map_ordering_tgt_alloc_from_to.c
@@ -0,0 +1,14 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <stdio.h>
+
+int main() {
+ int x = 111;
+#pragma omp target map(alloc : x) map(from : x) map(to : x) map(alloc : x)
+ {
+ printf("%d\n", x); // CHECK: 111
+ x = x + 111;
+ }
+
+ printf("%d\n", x); // CHECK: 222
+}
diff --git a/offload/test/mapping/map_ordering_tgt_alloc_present_tofrom.c b/offload/test/mapping/map_ordering_tgt_alloc_present_tofrom.c
new file mode 100644
index 0000000000000..f8f397efc4acd
--- /dev/null
+++ b/offload/test/mapping/map_ordering_tgt_alloc_present_tofrom.c
@@ -0,0 +1,27 @@
+// RUN: %libomptarget-compile-generic
+// RUN: %libomptarget-run-fail-generic 2>&1 \
+// RUN: | %fcheck-generic
+
+#include <stdio.h>
+
+int main() {
+ // CHECK: addr=0x[[#%x,HOST_ADDR:]], size=[[#%u,SIZE:]]
+ int x = 111;
+ fprintf(stderr, "addr=%p, size=%ld\n", &x, sizeof(x));
+// CHECK: omptarget message: device mapping required by 'present' map type
+// modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]]
+// bytes)
+// CHECK: omptarget error: Pointer 0x{{0*}}[[#HOST_ADDR]] was not present
+// on the device upon entry to the region.
+// ('present' map type modifier).
+// CHECK: omptarget error: Call to targetDataBegin failed, abort target.
+// CHECK: omptarget error: Failed to process data before launching the kernel.
+// CHECK: omptarget fatal error 1: failure of target construct while offloading
+// is mandatory
+#pragma omp target map(alloc : x) map(present, alloc : x) map(tofrom : x)
+ {
+ printf("%d\n", x);
+ }
+
+ return 0;
+}
diff --git a/offload/test/mapping/map_ordering_tgt_alloc_tofrom.c b/offload/test/mapping/map_ordering_tgt_alloc_tofrom.c
new file mode 100644
index 0000000000000..c76e2b4bafa1a
--- /dev/null
+++ b/offload/test/mapping/map_ordering_tgt_alloc_tofrom.c
@@ -0,0 +1,14 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <stdio.h>
+
+int main() {
+ int x = 111;
+#pragma omp target map(alloc : x) map(tofrom : x) map(alloc : x)
+ {
+ printf("%d\n", x); // CHECK: 111
+ x = x + 111;
+ }
+
+ printf("%d\n", x); // CHECK: 222
+}
diff --git a/offload/test/mapping/map_ordering_tgt_data_alloc_from.c b/offload/test/mapping/map_ordering_tgt_data_alloc_from.c
new file mode 100644
index 0000000000000..e5905460bea19
--- /dev/null
+++ b/offload/test/mapping/map_ordering_tgt_data_alloc_from.c
@@ -0,0 +1,14 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <stdio.h>
+
+int main() {
+ int x = 111;
+#pragma omp target data map(alloc : x) map(from : x) map(alloc : x)
+ {
+#pragma omp target map(present, alloc : x)
+ x = 222;
+ }
+
+ printf("%d\n", x); // CHECK: 222
+}
diff --git a/offload/test/mapping/map_ordering_tgt_data_alloc_to_from.c b/offload/test/mapping/map_ordering_tgt_data_alloc_to_from.c
new file mode 100644
index 0000000000000..1ed41200cecde
--- /dev/null
+++ b/offload/test/mapping/map_ordering_tgt_data_alloc_to_from.c
@@ -0,0 +1,17 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <stdio.h>
+
+int main() {
+ int x = 111;
+#pragma omp target data map(alloc : x) map(to : x) map(from : x) map(alloc : x)
+ {
+#pragma omp target map(present, alloc : x)
+ {
+ printf("%d\n", x); // CHECK: 111
+ x = x + 111;
+ }
+ }
+
+ printf("%d\n", x); // CHECK: 222
+}
diff --git a/offload/test/mapping/map_ordering_tgt_data_alloc_tofrom.c b/offload/test/mapping/map_ordering_tgt_data_alloc_tofrom.c
new file mode 100644
index 0000000000000..6db30d2aa7f9d
--- /dev/null
+++ b/offload/test/mapping/map_ordering_tgt_data_alloc_tofrom.c
@@ -0,0 +1,17 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <stdio.h>
+
+int main() {
+ int x = 111;
+#pragma omp target data map(alloc : x) map(tofrom : x) map(alloc : x)
+ {
+#pragma omp target map(present, alloc : x)
+ {
+ printf("%d\n", x); // CHECK: 111
+ x = x + 111;
+ }
+ }
+
+ printf("%d\n", x); // CHECK: 222
+}
>From b2470f4baea4bf612afbe4fcb43c1a279126092e Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Wed, 29 Oct 2025 15:35:57 -0700
Subject: [PATCH 02/16] Fix 'from' + 'delete', and multiple 'always,from'
entries.
---
offload/include/OpenMP/Mapping.h | 8 ++++
offload/libomptarget/omptarget.cpp | 47 +++++++++++++++----
...map_ordering_tgt_exit_data_always_always.c | 29 ++++++++++++
.../map_ordering_tgt_exit_data_delete_from.c | 21 +++++++++
4 files changed, 97 insertions(+), 8 deletions(-)
create mode 100644 offload/test/mapping/map_ordering_tgt_exit_data_always_always.c
create mode 100644 offload/test/mapping/map_ordering_tgt_exit_data_delete_from.c
diff --git a/offload/include/OpenMP/Mapping.h b/offload/include/OpenMP/Mapping.h
index 517f6c0a99244..686e72e1316f9 100644
--- a/offload/include/OpenMP/Mapping.h
+++ b/offload/include/OpenMP/Mapping.h
@@ -499,6 +499,14 @@ struct StateInfoTy {
/// occur due to the ref-count not being zero.
llvm::SmallSet<void *, 32> DeferredFromPtrs;
+ /// Host pointers for which we have attempted a FROM transfer at some point
+ /// during targetDataEnd. Used to avoid duplicate transfers.
+ llvm::SmallSet<void *, 32> TransferredFromPtrs;
+
+ /// Host pointers for which a DELETE entry was encountered, causing their
+ /// ref-count to have gone down to zero.
+ llvm::SmallSet<void *, 32> MarkedForDeletionPtrs;
+
StateInfoTy() = default;
// Delete copy constructor and copy assignment operator to prevent copying
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index bef1488b2956f..8ae6b6520ae9a 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -1086,6 +1086,7 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
HstPtrBegin, DataSize, UpdateRef, HasHoldModifier, !IsImplicit,
ForceDelete, /*FromDataEnd=*/true);
void *TgtPtrBegin = TPR.TargetPointer;
+
if (!TPR.isPresent() && !TPR.isHostPointer() &&
(DataSize || HasPresentModifier)) {
DP("Mapping does not exist (%s)\n",
@@ -1125,6 +1126,11 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
if (!TPR.isPresent())
continue;
+ // Track force-deleted pointers so we can use this information if we
+ // encounter FROM entries for the same pointer later on.
+ if (ForceDelete && TPR.Flags.IsLast)
+ StateInfo->MarkedForDeletionPtrs.insert(HstPtrBegin);
+
// Move data back to the host
const bool HasAlways = ArgTypes[I] & OMP_TGT_MAPTYPE_ALWAYS;
const bool HasFrom = ArgTypes[I] & OMP_TGT_MAPTYPE_FROM;
@@ -1141,15 +1147,40 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
return true;
};
+ // Lambda to check if this pointer was previously marked for deletion.
+ // Such a pointer would have had "IsLast" set to true when its DELETE entry
+ // was processed. So, the flag wouldn't be set for any FROM entries seen
+ // later on.
+ auto WasPreviouslyMarkedForDeletion = [&]() -> bool {
+ if (!StateInfo->MarkedForDeletionPtrs.contains(HstPtrBegin))
+ return false;
+ DP("Pointer HstPtr=" DPxMOD " was previously marked for deletion\n",
+ DPxPTR(HstPtrBegin));
+ return true;
+ };
+
+ bool FromCopyBackAlreadyDone =
+ StateInfo->TransferredFromPtrs.contains(HstPtrBegin);
bool IsMapFromOnNonHostNonZeroData =
HasFrom && !TPR.Flags.IsHostPointer && DataSize != 0;
- bool IsLastOrHasAlways = TPR.Flags.IsLast || HasAlways;
+ bool IsLastOrHasAlwaysOrWasForceDeleted =
+ TPR.Flags.IsLast || HasAlways || WasPreviouslyMarkedForDeletion();
+
+ if (!FromCopyBackAlreadyDone &&
+ ((IsMapFromOnNonHostNonZeroData &&
+ IsLastOrHasAlwaysOrWasForceDeleted) ||
+ // Even if are not looking at an entry with FROM map-type, if there
+ // were any previously deferred FROM transfers for this pointer, we
+ // should do them when the ref-count goes down to zero.
+ (TPR.Flags.IsLast && HasDeferredMapFrom()))) {
+ // Track that we're doing a FROM transfer for this pointer
+ // NOTE: If we don't care about the case of multiple different maps with
+ // from, always, or multiple map(from)s seen after a map(delete), e.g.
+ // ... map(always, from: x) map(always, from: x)
+ // ... map(delete: x) map(from: x) map(from: x)
+ // Then we can forego tacking TransferredFromPtrs.
+ StateInfo->TransferredFromPtrs.insert(HstPtrBegin);
- if ((IsMapFromOnNonHostNonZeroData && IsLastOrHasAlways) ||
- // Even if are not looking at an entry with FROM map-type, if there were
- // any previously deferred FROM transfers for this pointer, we should
- // do them when the ref-count goes down to zero.
- (TPR.Flags.IsLast && HasDeferredMapFrom())) {
DP("Moving %" PRId64 " bytes (tgt:" DPxMOD ") -> (hst:" DPxMOD ")\n",
DataSize, DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin));
TIMESCOPE_WITH_DETAILS_AND_IDENT(
@@ -1179,8 +1210,8 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
OFFLOAD_SUCCESS)
return OFFLOAD_FAIL;
}
- } else if (IsMapFromOnNonHostNonZeroData && !IsLastOrHasAlways &&
- !IsMemberOf) {
+ } else if (!FromCopyBackAlreadyDone && IsMapFromOnNonHostNonZeroData &&
+ !IsLastOrHasAlwaysOrWasForceDeleted && !IsMemberOf) {
// We can have cases like the following:
// map(alloc: p[0:1]) map(from: p[0:1])
//
diff --git a/offload/test/mapping/map_ordering_tgt_exit_data_always_always.c b/offload/test/mapping/map_ordering_tgt_exit_data_always_always.c
new file mode 100644
index 0000000000000..ea8e61befff9b
--- /dev/null
+++ b/offload/test/mapping/map_ordering_tgt_exit_data_always_always.c
@@ -0,0 +1,29 @@
+// RUN: %libomptarget-compile-generic
+// RUN: env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefix=DEBUG -check-prefix=CHECK
+// REQUIRES: libomptarget-debug
+
+// There should only be one "from" data-transfer, despite the two duplicate
+// maps.
+
+#include <stdio.h>
+
+int main() {
+ int x = 111;
+#pragma omp target data map(alloc : x)
+ {
+#pragma omp target enter data map(alloc : x) map(to : x)
+ {
+#pragma omp target map(present, alloc : x)
+ {
+ printf("%d\n", x); // CHECK-NOT: 111
+ x = 222;
+ }
+ }
+#pragma omp target exit data map(always, from : x) map(always, from : x)
+ // DEBUG: omptarget --> Moving 4 bytes (tgt:0x{{.*}}) -> (hst:0x{{.*}})
+ // DEBUG-NOT: omptarget --> Moving 4 bytes
+ }
+
+ printf("%d\n", x); // CHECK: 222
+}
diff --git a/offload/test/mapping/map_ordering_tgt_exit_data_delete_from.c b/offload/test/mapping/map_ordering_tgt_exit_data_delete_from.c
new file mode 100644
index 0000000000000..9b2f534556dc0
--- /dev/null
+++ b/offload/test/mapping/map_ordering_tgt_exit_data_delete_from.c
@@ -0,0 +1,21 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <stdio.h>
+
+int main() {
+ int x = 111;
+#pragma omp target data map(alloc : x)
+ {
+#pragma omp target enter data map(alloc : x) map(to : x)
+ {
+#pragma omp target map(present, alloc : x)
+ {
+ printf("%d\n", x); // CHECK-NOT: 111
+ x = 222;
+ }
+ }
+#pragma omp target exit data map(from : x) map(delete : x)
+ }
+
+ printf("%d\n", x); // CHECK: 222
+}
>From a44a8dc295e0a4ac11f35e4511f10a78f500b74e Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Tue, 4 Nov 2025 16:34:39 -0800
Subject: [PATCH 03/16] Handle cases like map(delete:p[:]) map(from:q[0:1]).
---
offload/include/OpenMP/Mapping.h | 3 +-
offload/libomptarget/OpenMP/Mapping.cpp | 20 ++++---
offload/libomptarget/omptarget.cpp | 53 ++++++++++---------
.../mapping/map_ordering_tgt_alloc_from_to.c | 13 ++++-
.../map_ordering_tgt_exit_data_delete_from.c | 5 +-
...ng_tgt_exit_data_delete_from_assumedsize.c | 37 +++++++++++++
...ng_tgt_exit_data_from_delete_assumedsize.c | 38 +++++++++++++
7 files changed, 134 insertions(+), 35 deletions(-)
create mode 100644 offload/test/mapping/map_ordering_tgt_exit_data_delete_from_assumedsize.c
create mode 100644 offload/test/mapping/map_ordering_tgt_exit_data_from_delete_assumedsize.c
diff --git a/offload/include/OpenMP/Mapping.h b/offload/include/OpenMP/Mapping.h
index 686e72e1316f9..b150854e6e97d 100644
--- a/offload/include/OpenMP/Mapping.h
+++ b/offload/include/OpenMP/Mapping.h
@@ -497,7 +497,8 @@ struct StateInfoTy {
/// Host pointers that had a FROM entry, but for which a data transfer didn't
/// occur due to the ref-count not being zero.
- llvm::SmallSet<void *, 32> DeferredFromPtrs;
+ /// Key: host pointer, Value: data size.
+ llvm::DenseMap<void *, int64_t> DeferredFromEntries;
/// Host pointers for which we have attempted a FROM transfer at some point
/// during targetDataEnd. Used to avoid duplicate transfers.
diff --git a/offload/libomptarget/OpenMP/Mapping.cpp b/offload/libomptarget/OpenMP/Mapping.cpp
index a3f634bc0a9eb..e316af1876f4d 100644
--- a/offload/libomptarget/OpenMP/Mapping.cpp
+++ b/offload/libomptarget/OpenMP/Mapping.cpp
@@ -323,19 +323,27 @@ TargetPointerResultTy MappingInfoTy::getTargetPointer(
if (ReleaseHDTTMap)
HDTTMap.destroy();
- // If the target pointer is valid, and we need to transfer data, issue the
- // data transfer.
- auto WasNewlyAllocatedOnCurrentConstruct = [&]() {
+ // Lambda to check if this pointer was newly allocated on the current region.
+ // This is needed to handle cases when the TO entry is encounter after an
+ // alloc entry for the same pointer, which increased the ref-count from 0 to 1,
+ // has already been encountered before. But because the ref-count was already 1
+ // when TO was encountered, it wouldn't incur a transfer. e.g.
+ // ... map(alloc: x) map(to: x).
+ auto WasNewlyAllocatedForCurrentRegion = [&]() {
if (!StateInfo)
return false;
- return StateInfo->NewAllocations.contains(HstPtrBegin);
+ bool IsNewlyAllocated = StateInfo->NewAllocations.contains(HstPtrBegin);
+ if (IsNewlyAllocated)
+ DP("HstPtrBegin " DPxMOD " was newly allocated for the current region\n",
+ DPxPTR(HstPtrBegin));
+ return IsNewlyAllocated;
};
// Even if this isn't a new entry, we still need to do a data-transfer if
- // the pointer was newly allocated previously on the same construct.
+ // the pointer was newly allocated on the current target region.
if (LR.TPR.TargetPointer && !LR.TPR.Flags.IsHostPointer && HasFlagTo &&
(LR.TPR.Flags.IsNewEntry || HasFlagAlways ||
- WasNewlyAllocatedOnCurrentConstruct()) &&
+ WasNewlyAllocatedForCurrentRegion()) &&
Size != 0) {
// If we have something like:
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index 8ae6b6520ae9a..5ddcb33e693d2 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -1135,15 +1135,20 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
const bool HasAlways = ArgTypes[I] & OMP_TGT_MAPTYPE_ALWAYS;
const bool HasFrom = ArgTypes[I] & OMP_TGT_MAPTYPE_FROM;
const bool IsMemberOf = ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF;
+ int64_t TransferSize = DataSize; // Size for FROM data-transfer.
+
// Lambda to check if there was a previously deferred FROM for this pointer
- // due to its ref-count not being zero.
+ // due to its ref-count not being zero. Updates TransferSize if found.
auto HasDeferredMapFrom = [&]() -> bool {
- if (!StateInfo->DeferredFromPtrs.contains(HstPtrBegin))
+ auto It = StateInfo->DeferredFromEntries.find(HstPtrBegin);
+ if (It == StateInfo->DeferredFromEntries.end())
return false;
- DP("Found previously deferred FROM transfer for HstPtr=" DPxMOD "\n",
- DPxPTR(HstPtrBegin));
- // Remove it so we don't look at it again
- StateInfo->DeferredFromPtrs.erase(HstPtrBegin);
+ DP("Found previously deferred FROM transfer for HstPtr=" DPxMOD
+ ", with size "
+ "%" PRId64 "\n",
+ DPxPTR(HstPtrBegin), It->second);
+ TransferSize = It->second;
+ StateInfo->DeferredFromEntries.erase(It);
return true;
};
@@ -1151,6 +1156,11 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
// Such a pointer would have had "IsLast" set to true when its DELETE entry
// was processed. So, the flag wouldn't be set for any FROM entries seen
// later on.
+ // This is needed to handle cases like the following:
+ // p1 = p2 = &x;
+ // ... map(delete: p1[:]) map(from: p2[0:1])
+ // The ref-count becomes zero before encountering the FROM entry, but we
+ // still need to do a transfer, if it went from non-zero to zero.
auto WasPreviouslyMarkedForDeletion = [&]() -> bool {
if (!StateInfo->MarkedForDeletionPtrs.contains(HstPtrBegin))
return false;
@@ -1169,7 +1179,7 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
if (!FromCopyBackAlreadyDone &&
((IsMapFromOnNonHostNonZeroData &&
IsLastOrHasAlwaysOrWasForceDeleted) ||
- // Even if are not looking at an entry with FROM map-type, if there
+ // Even if we're not looking at an entry with FROM map-type, if there
// were any previously deferred FROM transfers for this pointer, we
// should do them when the ref-count goes down to zero.
(TPR.Flags.IsLast && HasDeferredMapFrom()))) {
@@ -1182,9 +1192,9 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
StateInfo->TransferredFromPtrs.insert(HstPtrBegin);
DP("Moving %" PRId64 " bytes (tgt:" DPxMOD ") -> (hst:" DPxMOD ")\n",
- DataSize, DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin));
+ TransferSize, DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin));
TIMESCOPE_WITH_DETAILS_AND_IDENT(
- "DevToHost", "Size=" + std::to_string(DataSize) + "B", Loc);
+ "DevToHost", "Size=" + std::to_string(TransferSize) + "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) {
@@ -1193,8 +1203,8 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
}
}
- Ret = Device.retrieveData(HstPtrBegin, TgtPtrBegin, DataSize, AsyncInfo,
- TPR.getEntry());
+ Ret = Device.retrieveData(HstPtrBegin, TgtPtrBegin, TransferSize,
+ AsyncInfo, TPR.getEntry());
if (Ret != OFFLOAD_SUCCESS) {
REPORT("Copying data from device failed.\n");
return OFFLOAD_FAIL;
@@ -1213,24 +1223,19 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
} else if (!FromCopyBackAlreadyDone && IsMapFromOnNonHostNonZeroData &&
!IsLastOrHasAlwaysOrWasForceDeleted && !IsMemberOf) {
// We can have cases like the following:
- // map(alloc: p[0:1]) map(from: p[0:1])
+ // ... map(storage: p1[0:1]) map(from: p1[0:1])
//
- // For such cases, if we have different entries for the two maps, we
- // may not see the ref-count go down to zero when handling the From entry.
+ // where it's possible that when the FROM entry is processed, the
+ // ref count is not zero, so no data transfer happens for it. But
+ // the ref-count can go down to zero by the end of the directive
+ // in which case a transfer should happen.
//
- // So, we defer the FROM data-transfer until the ref-count goes down to
- // zero (if it does).
+ // So, we keep track of any skipped FROM data-transfers, in case
+ // the ref-count goes down to zero later on.
//
// This should be limited to non-member-of entries because for member-of,
// their ref-count should go down only once as part of the parent.
- //
- // Also, we don't need to worry about cases like:
- // map(alloc: p[0:10]) map(from: p[0:1])
- //
- // because that is not OpenMP 6.0 compliant, so we can just save the
- // pointer without saving the size, and assume that the size for the
- // "alloc" map will match that of "from".
- StateInfo->DeferredFromPtrs.insert(HstPtrBegin);
+ StateInfo->DeferredFromEntries[HstPtrBegin] = DataSize;
DP("Deferring FROM map transfer for HstPtr=" DPxMOD ", Size=%" PRId64
"\n",
DPxPTR(HstPtrBegin), DataSize);
diff --git a/offload/test/mapping/map_ordering_tgt_alloc_from_to.c b/offload/test/mapping/map_ordering_tgt_alloc_from_to.c
index 67c88e7238842..ed20e98dde512 100644
--- a/offload/test/mapping/map_ordering_tgt_alloc_from_to.c
+++ b/offload/test/mapping/map_ordering_tgt_alloc_from_to.c
@@ -1,9 +1,20 @@
-// RUN: %libomptarget-compile-run-and-check-generic
+// RUN: %libomptarget-compile-generic -fopenmp-version=60
+// RUN: env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefix=DEBUG -check-prefix=CHECK
+// REQUIRES: libomptarget-debug
#include <stdio.h>
+// Even if the "alloc" and "from" are encountered before the "to",
+// there should be a data-transfer from host to device, as the
+// ref-count goes from 0 to 1 at the entry of the target region.
+
int main() {
int x = 111;
+ // DEBUG: omptarget --> HstPtrBegin 0x[[#%x,HOST_ADDR:]] was newly allocated
+ // DEBUG-SAME: for the current region
+ // DEBUG: omptarget --> Moving {{.*}} bytes
+ // DEBUG-SAME: (hst:0x{{0*}}[[#HOST_ADDR]]) -> (tgt:0x{{.*}})
#pragma omp target map(alloc : x) map(from : x) map(to : x) map(alloc : x)
{
printf("%d\n", x); // CHECK: 111
diff --git a/offload/test/mapping/map_ordering_tgt_exit_data_delete_from.c b/offload/test/mapping/map_ordering_tgt_exit_data_delete_from.c
index 9b2f534556dc0..ad7db66890edb 100644
--- a/offload/test/mapping/map_ordering_tgt_exit_data_delete_from.c
+++ b/offload/test/mapping/map_ordering_tgt_exit_data_delete_from.c
@@ -14,8 +14,7 @@ int main() {
x = 222;
}
}
-#pragma omp target exit data map(from : x) map(delete : x)
+#pragma omp target exit data map(delete : x) map(from : x) map(delete : x)
+ printf("%d\n", x); // CHECK: 222
}
-
- printf("%d\n", x); // CHECK: 222
}
diff --git a/offload/test/mapping/map_ordering_tgt_exit_data_delete_from_assumedsize.c b/offload/test/mapping/map_ordering_tgt_exit_data_delete_from_assumedsize.c
new file mode 100644
index 0000000000000..b9fa2b99b56d9
--- /dev/null
+++ b/offload/test/mapping/map_ordering_tgt_exit_data_delete_from_assumedsize.c
@@ -0,0 +1,37 @@
+// RUN: %libomptarget-compile-generic -fopenmp-version=60
+// RUN: env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefix=DEBUG -check-prefix=CHECK
+// REQUIRES: libomptarget-debug
+
+// The from on target_exit_data should result in a data-transfer of 4 bytes,
+// even if when "from" is honored, the ref-count hasn't gone down to 0.
+// It will eventually go down to 0 as part of the same exit_data due to the
+// "delete" on it.
+// This is a case that cannot be handled at compile time because the list-items
+// are not related.
+
+#include <stdio.h>
+
+int main() {
+ int x[10];
+ int *p1x, *p2x;
+ p1x = p2x = &x[0];
+
+#pragma omp target data map(alloc : x)
+ {
+#pragma omp target enter data map(alloc : x) map(to : x)
+ {
+#pragma omp target map(present, alloc : x)
+ {
+ printf("%d\n", x[0]); // CHECK-NOT: 111
+ x[0] = 222;
+ }
+ }
+// DEBUG: omptarget --> Found previously deferred FROM transfer
+// DEBUG-SAME: for HstPtr=0x[[#%x,HOST_ADDR:]], with size [[#%u,SIZE:]]
+// DEBUG: omptarget --> Moving [[#SIZE]] bytes
+// DEBUG-SAME: (tgt:0x{{.*}}) -> (hst:0x{{0*}}[[#HOST_ADDR]])
+#pragma omp target exit data map(delete : p1x[ : ]) map(from : p2x[0])
+ printf("%d\n", x[0]); // CHECK: 222
+ }
+}
diff --git a/offload/test/mapping/map_ordering_tgt_exit_data_from_delete_assumedsize.c b/offload/test/mapping/map_ordering_tgt_exit_data_from_delete_assumedsize.c
new file mode 100644
index 0000000000000..af54ef0b183b6
--- /dev/null
+++ b/offload/test/mapping/map_ordering_tgt_exit_data_from_delete_assumedsize.c
@@ -0,0 +1,38 @@
+// RUN: %libomptarget-compile-generic -fopenmp-version=60
+// RUN: env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefix=DEBUG -check-prefix=CHECK
+// REQUIRES: libomptarget-debug
+
+#include <stdio.h>
+
+// The from on target_exit_data should result in a data-transfer of 4 bytes,
+// even if when "delete" is honored first, and by the time "from" is
+// encountered, the ref-count had already been 0 (i.e. it's not transitioning
+// from non-zero to zero).
+// This is a case that cannot be handled at compile time because the list-items
+// are not related.
+
+#include <stdio.h>
+int main() {
+ int x[10];
+ int *p1x, *p2x;
+ p1x = p2x = &x[0];
+
+#pragma omp target data map(alloc : x)
+ {
+#pragma omp target enter data map(alloc : x) map(to : x)
+ {
+#pragma omp target map(present, alloc : x)
+ {
+ printf("%d\n", x[0]); // CHECK-NOT: 111
+ x[0] = 222;
+ }
+ }
+ // DEBUG: omptarget --> Pointer HstPtr=0x[[#%x,HOST_ADDR:]]
+ // DEBUG-SAME: was previously marked for deletion
+ // DEBUG: omptarget --> Moving {{.*}} bytes
+ // DEBUG-SAME: (tgt:0x{{.*}}) -> (hst:0x{{0*}}[[#HOST_ADDR]])
+#pragma omp target exit data map(from : p2x[0]) map(delete : p1x[ : ])
+ printf("%d\n", x[0]); // CHECK: 222
+ }
+}
>From 85caebbac0c00c24ec1aff2aea7d8688ae80f778 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Tue, 4 Nov 2025 17:32:20 -0800
Subject: [PATCH 04/16] Clang-format
---
offload/libomptarget/OpenMP/Mapping.cpp | 8 ++++----
offload/test/mapping/map_ordering_tgt_alloc_from_to.c | 8 ++++----
2 files changed, 8 insertions(+), 8 deletions(-)
diff --git a/offload/libomptarget/OpenMP/Mapping.cpp b/offload/libomptarget/OpenMP/Mapping.cpp
index e316af1876f4d..2286c422d41f2 100644
--- a/offload/libomptarget/OpenMP/Mapping.cpp
+++ b/offload/libomptarget/OpenMP/Mapping.cpp
@@ -325,16 +325,16 @@ TargetPointerResultTy MappingInfoTy::getTargetPointer(
// Lambda to check if this pointer was newly allocated on the current region.
// This is needed to handle cases when the TO entry is encounter after an
- // alloc entry for the same pointer, which increased the ref-count from 0 to 1,
- // has already been encountered before. But because the ref-count was already 1
- // when TO was encountered, it wouldn't incur a transfer. e.g.
+ // alloc entry for the same pointer, which increased the ref-count from 0 to
+ // 1, has already been encountered before. But because the ref-count was
+ // already 1 when TO was encountered, it wouldn't incur a transfer. e.g.
// ... map(alloc: x) map(to: x).
auto WasNewlyAllocatedForCurrentRegion = [&]() {
if (!StateInfo)
return false;
bool IsNewlyAllocated = StateInfo->NewAllocations.contains(HstPtrBegin);
if (IsNewlyAllocated)
- DP("HstPtrBegin " DPxMOD " was newly allocated for the current region\n",
+ DP("HstPtrBegin " DPxMOD " was newly allocated for the current region\n",
DPxPTR(HstPtrBegin));
return IsNewlyAllocated;
};
diff --git a/offload/test/mapping/map_ordering_tgt_alloc_from_to.c b/offload/test/mapping/map_ordering_tgt_alloc_from_to.c
index ed20e98dde512..cdcb62ea0ba8e 100644
--- a/offload/test/mapping/map_ordering_tgt_alloc_from_to.c
+++ b/offload/test/mapping/map_ordering_tgt_alloc_from_to.c
@@ -11,10 +11,10 @@
int main() {
int x = 111;
- // DEBUG: omptarget --> HstPtrBegin 0x[[#%x,HOST_ADDR:]] was newly allocated
- // DEBUG-SAME: for the current region
- // DEBUG: omptarget --> Moving {{.*}} bytes
- // DEBUG-SAME: (hst:0x{{0*}}[[#HOST_ADDR]]) -> (tgt:0x{{.*}})
+ // DEBUG: omptarget --> HstPtrBegin 0x[[#%x,HOST_ADDR:]] was newly allocated
+ // DEBUG-SAME: for the current region
+ // DEBUG: omptarget --> Moving {{.*}} bytes
+ // DEBUG-SAME: (hst:0x{{0*}}[[#HOST_ADDR]]) -> (tgt:0x{{.*}})
#pragma omp target map(alloc : x) map(from : x) map(to : x) map(alloc : x)
{
printf("%d\n", x); // CHECK: 111
>From de64b9c77d7686bdf7f0a3a8b9403e52f8c9b9c6 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Thu, 6 Nov 2025 16:52:29 -0800
Subject: [PATCH 05/16] Add two mapper tests, avoid a rendant map lookup.
---
offload/libomptarget/omptarget.cpp | 9 +++--
...ring_ptee_tgt_alloc_mapper_alloc_from_to.c | 35 ++++++++++++++++++
..._alloc_tgt_mapper_present_delete_from_to.c | 37 +++++++++++++++++++
.../mapping/map_ordering_tgt_alloc_from_to.c | 2 +-
4 files changed, 78 insertions(+), 5 deletions(-)
create mode 100644 offload/test/mapping/map_ordering_ptee_tgt_alloc_mapper_alloc_from_to.c
create mode 100644 offload/test/mapping/map_ordering_ptee_tgt_data_alloc_tgt_mapper_present_delete_from_to.c
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index 5ddcb33e693d2..06753ffe29f38 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -1173,12 +1173,13 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
StateInfo->TransferredFromPtrs.contains(HstPtrBegin);
bool IsMapFromOnNonHostNonZeroData =
HasFrom && !TPR.Flags.IsHostPointer && DataSize != 0;
- bool IsLastOrHasAlwaysOrWasForceDeleted =
- TPR.Flags.IsLast || HasAlways || WasPreviouslyMarkedForDeletion();
+ auto IsLastOrHasAlwaysOrWasForceDeleted = [&]() {
+ return TPR.Flags.IsLast || HasAlways || WasPreviouslyMarkedForDeletion();
+ };
if (!FromCopyBackAlreadyDone &&
((IsMapFromOnNonHostNonZeroData &&
- IsLastOrHasAlwaysOrWasForceDeleted) ||
+ IsLastOrHasAlwaysOrWasForceDeleted()) ||
// Even if we're not looking at an entry with FROM map-type, if there
// were any previously deferred FROM transfers for this pointer, we
// should do them when the ref-count goes down to zero.
@@ -1221,7 +1222,7 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
return OFFLOAD_FAIL;
}
} else if (!FromCopyBackAlreadyDone && IsMapFromOnNonHostNonZeroData &&
- !IsLastOrHasAlwaysOrWasForceDeleted && !IsMemberOf) {
+ !IsLastOrHasAlwaysOrWasForceDeleted() && !IsMemberOf) {
// We can have cases like the following:
// ... map(storage: p1[0:1]) map(from: p1[0:1])
//
diff --git a/offload/test/mapping/map_ordering_ptee_tgt_alloc_mapper_alloc_from_to.c b/offload/test/mapping/map_ordering_ptee_tgt_alloc_mapper_alloc_from_to.c
new file mode 100644
index 0000000000000..04b41265c0852
--- /dev/null
+++ b/offload/test/mapping/map_ordering_ptee_tgt_alloc_mapper_alloc_from_to.c
@@ -0,0 +1,35 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+// Since the allocation of the pointee happens as on the "target" construct (1),
+// the "to" transfer requested as part of the mapper should also happen.
+//
+// Similarly, the "from" transfer should also happen at the end of the target
+// construct, even if the ref-count of the pointee x has not gone down to 0
+// when "from" is encountered.
+
+// This currently fails, but should start passing once ATTACH-style maps are
+// enabled for mappers (#166874).
+// XFAIL: *
+
+#include <stdio.h>
+
+typedef struct {
+ int *p;
+ int *q;
+} S;
+#pragma omp declare mapper(my_mapper: S s) map(alloc: s.p, s.p[0:10]) map(from: s.p[0:10]) map(to: s.p[0:10]) map(alloc: s.p[0:10])
+
+ S s1;
+int main() {
+ int x[10];
+ x[1] = 111;
+ s1.q = s1.p = &x[0];
+
+ #pragma omp target map(alloc: s1.p[0:10]) map(mapper(my_mapper), tofrom: s1) // (1)
+ {
+ printf("%d\n", s1.p[1]); // CHECK: 111
+ s1.p[1] = s1.p[1] + 111;
+ }
+
+ printf("%d\n", s1.p[1]); // CHECK: 222
+}
diff --git a/offload/test/mapping/map_ordering_ptee_tgt_data_alloc_tgt_mapper_present_delete_from_to.c b/offload/test/mapping/map_ordering_ptee_tgt_data_alloc_tgt_mapper_present_delete_from_to.c
new file mode 100644
index 0000000000000..6bd8588f0b291
--- /dev/null
+++ b/offload/test/mapping/map_ordering_ptee_tgt_data_alloc_tgt_mapper_present_delete_from_to.c
@@ -0,0 +1,37 @@
+// RUN: %libomptarget-compile-generic -fopenmp-version=60
+// RUN: %libomptarget-run-generic | %fcheck-generic
+
+// The "present" check should pass on the "target" construct // (2),
+// and there should be no "to" transfer, because the pointee "x" is already
+// present (because of (1)).
+// However, there should be a "from" transfer at the end of (2) because of the
+// "delete" on the mapper.
+
+// This currently fails, but should start passing once ATTACH-style maps are
+// enabled for mappers (#166874).
+// XFAIL: *
+
+#include <stdio.h>
+
+typedef struct {
+ int *p;
+ int *q;
+} S;
+#pragma omp declare mapper(my_mapper: S s) map(alloc: s.p) map(alloc, present: s.p[0:10]) map(delete: s.q[:]) map(from: s.p[0:10]) map(to: s.p[0:10]) map(alloc: s.p[0:10])
+
+S s1;
+int main() {
+ int x[10];
+ x[1] = 111;
+ s1.q = s1.p = &x[0];
+
+ #pragma omp target data map(alloc: x) // (1)
+ {
+ #pragma omp target map(mapper(my_mapper), tofrom: s1) // (2)
+ {
+ printf("%d\n", s1.p[1]); // CHECK-NOT: 111
+ s1.p[1] = 222;
+ }
+ printf("%d\n", s1.p[1]); // CHECK: 222
+ }
+}
diff --git a/offload/test/mapping/map_ordering_tgt_alloc_from_to.c b/offload/test/mapping/map_ordering_tgt_alloc_from_to.c
index cdcb62ea0ba8e..71f68134c7317 100644
--- a/offload/test/mapping/map_ordering_tgt_alloc_from_to.c
+++ b/offload/test/mapping/map_ordering_tgt_alloc_from_to.c
@@ -1,4 +1,4 @@
-// RUN: %libomptarget-compile-generic -fopenmp-version=60
+// RUN: %libomptarget-compile-generic
// RUN: env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 \
// RUN: | %fcheck-generic -check-prefix=DEBUG -check-prefix=CHECK
// REQUIRES: libomptarget-debug
>From 6eeec697fdc2270e0cc8f7b78ca1f0faffb2d1dc Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Thu, 6 Nov 2025 16:59:21 -0800
Subject: [PATCH 06/16] Change Deferred to Skipped, since that's more accurate
---
offload/include/OpenMP/Mapping.h | 12 ++++++------
offload/libomptarget/omptarget.cpp | 20 ++++++++++----------
2 files changed, 16 insertions(+), 16 deletions(-)
diff --git a/offload/include/OpenMP/Mapping.h b/offload/include/OpenMP/Mapping.h
index b150854e6e97d..0ec85b9dea344 100644
--- a/offload/include/OpenMP/Mapping.h
+++ b/offload/include/OpenMP/Mapping.h
@@ -484,9 +484,9 @@ struct AttachMapInfo {
MapType(Type), Pointername(Name) {}
};
-/// Structure to track new allocations, ATTACH entries and deferred data
-/// transfer information for a given construct, across recursive calls (for
-/// handling mappers) to targetDataBegin/targetDataEnd.
+/// Structure to track new allocations, ATTACH entries, DELETE entries and
+/// skipped FROM data transfer information for a given construct, across
+/// recursive calls (for handling mappers) to targetDataBegin/targetDataEnd.
struct StateInfoTy {
/// ATTACH map entries for deferred processing until all other maps are done.
llvm::SmallVector<AttachMapInfo> AttachEntries;
@@ -495,10 +495,10 @@ struct StateInfoTy {
/// Key: host pointer, Value: allocation size.
llvm::DenseMap<void *, int64_t> NewAllocations;
- /// Host pointers that had a FROM entry, but for which a data transfer didn't
- /// occur due to the ref-count not being zero.
+ /// Host pointers that had a FROM entry, but for which a data transfer was
+ /// skipped due to the ref-count not being zero.
/// Key: host pointer, Value: data size.
- llvm::DenseMap<void *, int64_t> DeferredFromEntries;
+ llvm::DenseMap<void *, int64_t> SkippedFromEntries;
/// Host pointers for which we have attempted a FROM transfer at some point
/// during targetDataEnd. Used to avoid duplicate transfers.
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index 06753ffe29f38..03ed06be3f188 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -1137,18 +1137,18 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
const bool IsMemberOf = ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF;
int64_t TransferSize = DataSize; // Size for FROM data-transfer.
- // Lambda to check if there was a previously deferred FROM for this pointer
+ // Lambda to check if there was a previously skipped FROM for this pointer
// due to its ref-count not being zero. Updates TransferSize if found.
- auto HasDeferredMapFrom = [&]() -> bool {
- auto It = StateInfo->DeferredFromEntries.find(HstPtrBegin);
- if (It == StateInfo->DeferredFromEntries.end())
+ auto HasSkippedMapFrom = [&]() -> bool {
+ auto It = StateInfo->SkippedFromEntries.find(HstPtrBegin);
+ if (It == StateInfo->SkippedFromEntries.end())
return false;
- DP("Found previously deferred FROM transfer for HstPtr=" DPxMOD
+ DP("Found previously skipped FROM transfer for HstPtr=" DPxMOD
", with size "
"%" PRId64 "\n",
DPxPTR(HstPtrBegin), It->second);
TransferSize = It->second;
- StateInfo->DeferredFromEntries.erase(It);
+ StateInfo->SkippedFromEntries.erase(It);
return true;
};
@@ -1181,9 +1181,9 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
((IsMapFromOnNonHostNonZeroData &&
IsLastOrHasAlwaysOrWasForceDeleted()) ||
// Even if we're not looking at an entry with FROM map-type, if there
- // were any previously deferred FROM transfers for this pointer, we
+ // were any previously skipped FROM transfers for this pointer, we
// should do them when the ref-count goes down to zero.
- (TPR.Flags.IsLast && HasDeferredMapFrom()))) {
+ (TPR.Flags.IsLast && HasSkippedMapFrom()))) {
// Track that we're doing a FROM transfer for this pointer
// NOTE: If we don't care about the case of multiple different maps with
// from, always, or multiple map(from)s seen after a map(delete), e.g.
@@ -1236,8 +1236,8 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
//
// This should be limited to non-member-of entries because for member-of,
// their ref-count should go down only once as part of the parent.
- StateInfo->DeferredFromEntries[HstPtrBegin] = DataSize;
- DP("Deferring FROM map transfer for HstPtr=" DPxMOD ", Size=%" PRId64
+ StateInfo->SkippedFromEntries[HstPtrBegin] = DataSize;
+ DP("Skipping FROM map transfer for HstPtr=" DPxMOD ", Size=%" PRId64
"\n",
DPxPTR(HstPtrBegin), DataSize);
}
>From 151dc09ae58b5e864d5d9b03b6a78f6dab1dda3a Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Thu, 6 Nov 2025 17:00:02 -0800
Subject: [PATCH 07/16] Clang-format
---
offload/libomptarget/omptarget.cpp | 3 +--
.../map_ordering_ptee_tgt_alloc_mapper_alloc_from_to.c | 8 +++++---
...tee_tgt_data_alloc_tgt_mapper_present_delete_from_to.c | 8 +++++---
3 files changed, 11 insertions(+), 8 deletions(-)
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index 03ed06be3f188..e5161699ad337 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -1237,8 +1237,7 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
// This should be limited to non-member-of entries because for member-of,
// their ref-count should go down only once as part of the parent.
StateInfo->SkippedFromEntries[HstPtrBegin] = DataSize;
- DP("Skipping FROM map transfer for HstPtr=" DPxMOD ", Size=%" PRId64
- "\n",
+ DP("Skipping FROM map transfer for HstPtr=" DPxMOD ", Size=%" PRId64 "\n",
DPxPTR(HstPtrBegin), DataSize);
}
diff --git a/offload/test/mapping/map_ordering_ptee_tgt_alloc_mapper_alloc_from_to.c b/offload/test/mapping/map_ordering_ptee_tgt_alloc_mapper_alloc_from_to.c
index 04b41265c0852..3ad8748cb0572 100644
--- a/offload/test/mapping/map_ordering_ptee_tgt_alloc_mapper_alloc_from_to.c
+++ b/offload/test/mapping/map_ordering_ptee_tgt_alloc_mapper_alloc_from_to.c
@@ -17,15 +17,17 @@ typedef struct {
int *p;
int *q;
} S;
-#pragma omp declare mapper(my_mapper: S s) map(alloc: s.p, s.p[0:10]) map(from: s.p[0:10]) map(to: s.p[0:10]) map(alloc: s.p[0:10])
+#pragma omp declare mapper(my_mapper : S s) map(alloc : s.p, s.p[0 : 10]) \
+ map(from : s.p[0 : 10]) map(to : s.p[0 : 10]) map(alloc : s.p[0 : 10])
- S s1;
+S s1;
int main() {
int x[10];
x[1] = 111;
s1.q = s1.p = &x[0];
- #pragma omp target map(alloc: s1.p[0:10]) map(mapper(my_mapper), tofrom: s1) // (1)
+#pragma omp target map(alloc : s1.p[0 : 10]) \
+ map(mapper(my_mapper), tofrom : s1) // (1)
{
printf("%d\n", s1.p[1]); // CHECK: 111
s1.p[1] = s1.p[1] + 111;
diff --git a/offload/test/mapping/map_ordering_ptee_tgt_data_alloc_tgt_mapper_present_delete_from_to.c b/offload/test/mapping/map_ordering_ptee_tgt_data_alloc_tgt_mapper_present_delete_from_to.c
index 6bd8588f0b291..5fb196e31d9c2 100644
--- a/offload/test/mapping/map_ordering_ptee_tgt_data_alloc_tgt_mapper_present_delete_from_to.c
+++ b/offload/test/mapping/map_ordering_ptee_tgt_data_alloc_tgt_mapper_present_delete_from_to.c
@@ -17,7 +17,9 @@ typedef struct {
int *p;
int *q;
} S;
-#pragma omp declare mapper(my_mapper: S s) map(alloc: s.p) map(alloc, present: s.p[0:10]) map(delete: s.q[:]) map(from: s.p[0:10]) map(to: s.p[0:10]) map(alloc: s.p[0:10])
+#pragma omp declare mapper(my_mapper : S s) map(alloc : s.p) \
+ map(alloc, present : s.p[0 : 10]) map(delete : s.q[ : ]) \
+ map(from : s.p[0 : 10]) map(to : s.p[0 : 10]) map(alloc : s.p[0 : 10])
S s1;
int main() {
@@ -25,9 +27,9 @@ int main() {
x[1] = 111;
s1.q = s1.p = &x[0];
- #pragma omp target data map(alloc: x) // (1)
+#pragma omp target data map(alloc : x) // (1)
{
- #pragma omp target map(mapper(my_mapper), tofrom: s1) // (2)
+#pragma omp target map(mapper(my_mapper), tofrom : s1) // (2)
{
printf("%d\n", s1.p[1]); // CHECK-NOT: 111
s1.p[1] = 222;
>From bd71a220a50c9a04ed50bda091ce419b388b611c Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Thu, 6 Nov 2025 17:07:40 -0800
Subject: [PATCH 08/16] Fix some tests.
---
.../test/mapping/map_ordering_tgt_exit_data_always_always.c | 2 +-
.../map_ordering_tgt_exit_data_delete_from_assumedsize.c | 4 ++--
.../map_ordering_tgt_exit_data_from_delete_assumedsize.c | 2 +-
3 files changed, 4 insertions(+), 4 deletions(-)
diff --git a/offload/test/mapping/map_ordering_tgt_exit_data_always_always.c b/offload/test/mapping/map_ordering_tgt_exit_data_always_always.c
index ea8e61befff9b..43f3aae26b009 100644
--- a/offload/test/mapping/map_ordering_tgt_exit_data_always_always.c
+++ b/offload/test/mapping/map_ordering_tgt_exit_data_always_always.c
@@ -16,7 +16,7 @@ int main() {
{
#pragma omp target map(present, alloc : x)
{
- printf("%d\n", x); // CHECK-NOT: 111
+ printf("In tgt: %d\n", x); // CHECK-NOT: In tgt: 111
x = 222;
}
}
diff --git a/offload/test/mapping/map_ordering_tgt_exit_data_delete_from_assumedsize.c b/offload/test/mapping/map_ordering_tgt_exit_data_delete_from_assumedsize.c
index b9fa2b99b56d9..f8a117fc0ecf7 100644
--- a/offload/test/mapping/map_ordering_tgt_exit_data_delete_from_assumedsize.c
+++ b/offload/test/mapping/map_ordering_tgt_exit_data_delete_from_assumedsize.c
@@ -23,11 +23,11 @@ int main() {
{
#pragma omp target map(present, alloc : x)
{
- printf("%d\n", x[0]); // CHECK-NOT: 111
+ printf("In tgt: %d\n", x[0]); // CHECK-NOT: In tgt: 111
x[0] = 222;
}
}
-// DEBUG: omptarget --> Found previously deferred FROM transfer
+// DEBUG: omptarget --> Found previously skipped FROM transfer
// DEBUG-SAME: for HstPtr=0x[[#%x,HOST_ADDR:]], with size [[#%u,SIZE:]]
// DEBUG: omptarget --> Moving [[#SIZE]] bytes
// DEBUG-SAME: (tgt:0x{{.*}}) -> (hst:0x{{0*}}[[#HOST_ADDR]])
diff --git a/offload/test/mapping/map_ordering_tgt_exit_data_from_delete_assumedsize.c b/offload/test/mapping/map_ordering_tgt_exit_data_from_delete_assumedsize.c
index af54ef0b183b6..3f5b08d3473f8 100644
--- a/offload/test/mapping/map_ordering_tgt_exit_data_from_delete_assumedsize.c
+++ b/offload/test/mapping/map_ordering_tgt_exit_data_from_delete_assumedsize.c
@@ -24,7 +24,7 @@ int main() {
{
#pragma omp target map(present, alloc : x)
{
- printf("%d\n", x[0]); // CHECK-NOT: 111
+ printf("In tgt: %d\n", x[0]); // CHECK-NOT: In tgt: 111
x[0] = 222;
}
}
>From ab3ea51e12c4022dffd350770613977f580d31de Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Wed, 12 Nov 2025 15:56:01 -0800
Subject: [PATCH 09/16] Support case of FROM pointer falling in the middle of
the last entry.
---
offload/include/OpenMP/Mapping.h | 6 +-
offload/libomptarget/omptarget.cpp | 157 +++++++++++++++++++----------
2 files changed, 105 insertions(+), 58 deletions(-)
diff --git a/offload/include/OpenMP/Mapping.h b/offload/include/OpenMP/Mapping.h
index 0ec85b9dea344..49e0b8d5ad4f0 100644
--- a/offload/include/OpenMP/Mapping.h
+++ b/offload/include/OpenMP/Mapping.h
@@ -504,9 +504,9 @@ struct StateInfoTy {
/// during targetDataEnd. Used to avoid duplicate transfers.
llvm::SmallSet<void *, 32> TransferredFromPtrs;
- /// Host pointers for which a DELETE entry was encountered, causing their
- /// ref-count to have gone down to zero.
- llvm::SmallSet<void *, 32> MarkedForDeletionPtrs;
+ /// Starting host address and size of the DELETE entries previouly processed
+ /// for the current region. Key: host pointer, Value: allocated size.
+ llvm::DenseMap<void *, int64_t> DeleteEntries;
StateInfoTy() = default;
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index e5161699ad337..60a5ff37d4caf 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -1128,8 +1128,17 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
// Track force-deleted pointers so we can use this information if we
// encounter FROM entries for the same pointer later on.
- if (ForceDelete && TPR.Flags.IsLast)
- StateInfo->MarkedForDeletionPtrs.insert(HstPtrBegin);
+ if (ForceDelete && TPR.Flags.IsLast) {
+ // For assumed-size arrays like map(delete: p[:]), the compiler provides
+ // no size information, so we need to get the actual allocated extent from
+ // the HDTT entry.
+ int64_t AllocatedSize =
+ TPR.getEntry()->HstPtrEnd - TPR.getEntry()->HstPtrBegin;
+ DP("Marking HstPtr=" DPxMOD " for deletion with allocated size=%" PRId64
+ "\n",
+ DPxPTR(HstPtrBegin), AllocatedSize);
+ StateInfo->DeleteEntries[HstPtrBegin] = AllocatedSize;
+ }
// Move data back to the host
const bool HasAlways = ArgTypes[I] & OMP_TGT_MAPTYPE_ALWAYS;
@@ -1137,19 +1146,38 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
const bool IsMemberOf = ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF;
int64_t TransferSize = DataSize; // Size for FROM data-transfer.
- // Lambda to check if there was a previously skipped FROM for this pointer
- // due to its ref-count not being zero. Updates TransferSize if found.
- auto HasSkippedMapFrom = [&]() -> bool {
- auto It = StateInfo->SkippedFromEntries.find(HstPtrBegin);
- if (It == StateInfo->SkippedFromEntries.end())
- return false;
- DP("Found previously skipped FROM transfer for HstPtr=" DPxMOD
- ", with size "
- "%" PRId64 "\n",
- DPxPTR(HstPtrBegin), It->second);
- TransferSize = It->second;
- StateInfo->SkippedFromEntries.erase(It);
- return true;
+ // Lambda to perform the actual FROM data retrieval from device to host
+ auto PerformFromRetrieval = [&](void *HstPtr, void *TgtPtr, int64_t Size,
+ HostDataToTargetTy *Entry) -> int {
+ DP("Moving %" PRId64 " bytes (tgt:" DPxMOD ") -> (hst:" DPxMOD ")\n",
+ Size, DPxPTR(TgtPtr), DPxPTR(HstPtr));
+ TIMESCOPE_WITH_DETAILS_AND_IDENT(
+ "DevToHost", "Size=" + std::to_string(Size) + "B", Loc);
+ // Wait for any previous transfer if an event is present.
+ if (void *Event = Entry->getEvent()) {
+ if (Device.waitEvent(Event, AsyncInfo) != OFFLOAD_SUCCESS) {
+ REPORT("Failed to wait for event " DPxMOD ".\n", DPxPTR(Event));
+ return OFFLOAD_FAIL;
+ }
+ }
+
+ int Ret = Device.retrieveData(HstPtr, TgtPtr, Size, AsyncInfo, Entry);
+ if (Ret != OFFLOAD_SUCCESS) {
+ REPORT("Copying data from device failed.\n");
+ return OFFLOAD_FAIL;
+ }
+
+ // As we are expecting to delete the entry the d2h copy might race
+ // with another one that also tries to delete the entry. This happens
+ // as the entry can be reused and the reuse might happen after the
+ // copy-back was issued but before it completed. Since the reuse might
+ // also copy-back a value we would race.
+ if (TPR.Flags.IsLast) {
+ if (Entry->addEventIfNecessary(Device, AsyncInfo) != OFFLOAD_SUCCESS)
+ return OFFLOAD_FAIL;
+ }
+
+ return OFFLOAD_SUCCESS;
};
// Lambda to check if this pointer was previously marked for deletion.
@@ -1162,11 +1190,21 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
// The ref-count becomes zero before encountering the FROM entry, but we
// still need to do a transfer, if it went from non-zero to zero.
auto WasPreviouslyMarkedForDeletion = [&]() -> bool {
- if (!StateInfo->MarkedForDeletionPtrs.contains(HstPtrBegin))
- return false;
- DP("Pointer HstPtr=" DPxMOD " was previously marked for deletion\n",
- DPxPTR(HstPtrBegin));
- return true;
+ // Check if this pointer falls within the range of any deleted entry
+ for (const auto &DeleteEntry : StateInfo->DeleteEntries) {
+ void *DeletePtr = DeleteEntry.first;
+ int64_t DeleteSize = DeleteEntry.second;
+ if (HstPtrBegin >= DeletePtr &&
+ HstPtrBegin < (void *)((char *)DeletePtr + DeleteSize)) {
+ DP("Pointer HstPtr=" DPxMOD
+ " falls within a range previously marked for deletion [" DPxMOD
+ ", " DPxMOD ") with size=%" PRId64 "\n",
+ DPxPTR(HstPtrBegin), DPxPTR(DeletePtr),
+ DPxPTR((char *)DeletePtr + DeleteSize), DeleteSize);
+ return true;
+ }
+ }
+ return false;
};
bool FromCopyBackAlreadyDone =
@@ -1177,13 +1215,8 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
return TPR.Flags.IsLast || HasAlways || WasPreviouslyMarkedForDeletion();
};
- if (!FromCopyBackAlreadyDone &&
- ((IsMapFromOnNonHostNonZeroData &&
- IsLastOrHasAlwaysOrWasForceDeleted()) ||
- // Even if we're not looking at an entry with FROM map-type, if there
- // were any previously skipped FROM transfers for this pointer, we
- // should do them when the ref-count goes down to zero.
- (TPR.Flags.IsLast && HasSkippedMapFrom()))) {
+ if (!FromCopyBackAlreadyDone && (IsMapFromOnNonHostNonZeroData &&
+ IsLastOrHasAlwaysOrWasForceDeleted())) {
// Track that we're doing a FROM transfer for this pointer
// NOTE: If we don't care about the case of multiple different maps with
// from, always, or multiple map(from)s seen after a map(delete), e.g.
@@ -1192,35 +1225,10 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
// Then we can forego tacking TransferredFromPtrs.
StateInfo->TransferredFromPtrs.insert(HstPtrBegin);
- DP("Moving %" PRId64 " bytes (tgt:" DPxMOD ") -> (hst:" DPxMOD ")\n",
- TransferSize, DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin));
- TIMESCOPE_WITH_DETAILS_AND_IDENT(
- "DevToHost", "Size=" + std::to_string(TransferSize) + "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));
- return OFFLOAD_FAIL;
- }
- }
-
- Ret = Device.retrieveData(HstPtrBegin, TgtPtrBegin, TransferSize,
- AsyncInfo, TPR.getEntry());
- if (Ret != OFFLOAD_SUCCESS) {
- REPORT("Copying data from device failed.\n");
+ Ret = PerformFromRetrieval(HstPtrBegin, TgtPtrBegin, TransferSize,
+ TPR.getEntry());
+ if (Ret != OFFLOAD_SUCCESS)
return OFFLOAD_FAIL;
- }
-
- // As we are expecting to delete the entry the d2h copy might race
- // with another one that also tries to delete the entry. This happens
- // as the entry can be reused and the reuse might happen after the
- // copy-back was issued but before it completed. Since the reuse might
- // also copy-back a value we would race.
- if (TPR.Flags.IsLast) {
- if (TPR.getEntry()->addEventIfNecessary(Device, AsyncInfo) !=
- OFFLOAD_SUCCESS)
- return OFFLOAD_FAIL;
- }
} else if (!FromCopyBackAlreadyDone && IsMapFromOnNonHostNonZeroData &&
!IsLastOrHasAlwaysOrWasForceDeleted() && !IsMemberOf) {
// We can have cases like the following:
@@ -1239,6 +1247,45 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
StateInfo->SkippedFromEntries[HstPtrBegin] = DataSize;
DP("Skipping FROM map transfer for HstPtr=" DPxMOD ", Size=%" PRId64 "\n",
DPxPTR(HstPtrBegin), DataSize);
+ } else if (!FromCopyBackAlreadyDone && TPR.Flags.IsLast) {
+ // Even if this is not a FROM or DELETE entry, if the ref-count went to
+ // zero (IsLast=true), we should perform any previously skipped FROM
+ // transfers that fall within this entry's range.
+ SmallVector<void *, 32> ToRemove;
+
+ for (auto &SkippedEntry : StateInfo->SkippedFromEntries) {
+ void *SkippedPtr = SkippedEntry.first;
+ int64_t SkippedSize = SkippedEntry.second;
+ uintptr_t SkippedBegin = (uintptr_t)SkippedPtr;
+
+ uintptr_t EntryBegin = TPR.getEntry()->HstPtrBegin;
+ uintptr_t EntryEnd = TPR.getEntry()->HstPtrEnd;
+
+ // Check if skipped entry overlaps or is contained within current entry
+ if (SkippedBegin >= EntryBegin && SkippedBegin < EntryEnd) {
+ DP("Found skipped FROM within deleted region: HstPtr=" DPxMOD
+ " size=%" PRId64 " within [" DPxMOD ", " DPxMOD ")\n",
+ DPxPTR(SkippedPtr), SkippedSize, DPxPTR(EntryBegin),
+ DPxPTR(EntryEnd));
+
+ // Calculate offset within the target pointer
+ int64_t Offset = SkippedBegin - EntryBegin;
+ void *SkippedTgtPtr = (void *)((char *)TgtPtrBegin + Offset);
+
+ // Perform the retrieval for this skipped entry
+ int Ret = PerformFromRetrieval((void *)SkippedBegin, SkippedTgtPtr,
+ SkippedSize, TPR.getEntry());
+ if (Ret != OFFLOAD_SUCCESS)
+ return OFFLOAD_FAIL;
+
+ StateInfo->TransferredFromPtrs.insert(SkippedPtr);
+ ToRemove.push_back(SkippedPtr);
+ }
+ }
+
+ // Remove processed entries
+ for (void *Ptr : ToRemove)
+ StateInfo->SkippedFromEntries.erase(Ptr);
}
// Add pointer to the buffer for post-synchronize processing.
>From c77cd47171f9aaa959ed1b49eacca0dad2b2066f Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Wed, 12 Nov 2025 16:41:02 -0800
Subject: [PATCH 10/16] Support TO/PRESENT where the begin address falls in the
middle of a prior new allocation.
---
offload/include/OpenMP/Mapping.h | 13 +++++++++++++
offload/libomptarget/OpenMP/Mapping.cpp | 9 +++++----
offload/libomptarget/omptarget.cpp | 26 ++++++++-----------------
3 files changed, 26 insertions(+), 22 deletions(-)
diff --git a/offload/include/OpenMP/Mapping.h b/offload/include/OpenMP/Mapping.h
index 49e0b8d5ad4f0..b20739093030a 100644
--- a/offload/include/OpenMP/Mapping.h
+++ b/offload/include/OpenMP/Mapping.h
@@ -513,6 +513,19 @@ struct StateInfoTy {
// Delete copy constructor and copy assignment operator to prevent copying
StateInfoTy(const StateInfoTy &) = delete;
StateInfoTy &operator=(const StateInfoTy &) = delete;
+
+ /// Check if a pointer falls within any of the newly allocated ranges.
+ /// Returns true if the pointer is within a newly allocated region.
+ bool wasNewlyAllocated(void *Ptr) const {
+ return std::any_of(
+ NewAllocations.begin(), NewAllocations.end(), [&](const auto &Alloc) {
+ void *AllocPtr = Alloc.first;
+ int64_t AllocSize = Alloc.second;
+ return Ptr >= AllocPtr &&
+ Ptr < reinterpret_cast<void *>(
+ reinterpret_cast<char *>(AllocPtr) + AllocSize);
+ });
+ }
};
// Function pointer type for targetData* functions (targetDataBegin,
diff --git a/offload/libomptarget/OpenMP/Mapping.cpp b/offload/libomptarget/OpenMP/Mapping.cpp
index 2286c422d41f2..3cfe0c57f0d44 100644
--- a/offload/libomptarget/OpenMP/Mapping.cpp
+++ b/offload/libomptarget/OpenMP/Mapping.cpp
@@ -328,15 +328,16 @@ TargetPointerResultTy MappingInfoTy::getTargetPointer(
// alloc entry for the same pointer, which increased the ref-count from 0 to
// 1, has already been encountered before. But because the ref-count was
// already 1 when TO was encountered, it wouldn't incur a transfer. e.g.
- // ... map(alloc: x) map(to: x).
+ // int *xp = &x[0];
+ // ... map(alloc: x[:]) map(to: xp[1]).
auto WasNewlyAllocatedForCurrentRegion = [&]() {
if (!StateInfo)
return false;
- bool IsNewlyAllocated = StateInfo->NewAllocations.contains(HstPtrBegin);
- if (IsNewlyAllocated)
+ bool WasNewlyAllocated = StateInfo->wasNewlyAllocated(HstPtrBegin);
+ if (WasNewlyAllocated)
DP("HstPtrBegin " DPxMOD " was newly allocated for the current region\n",
DPxPTR(HstPtrBegin));
- return IsNewlyAllocated;
+ return WasNewlyAllocated;
};
// Even if this isn't a new entry, we still need to do a data-transfer if
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index 60a5ff37d4caf..0cf85b5d990b8 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -666,17 +666,14 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
: "device failure or illegal mapping");
return OFFLOAD_FAIL;
} else if (TgtPtrBegin && HasPresentModifier &&
- StateInfo->NewAllocations.contains(HstPtrBegin)) {
+ StateInfo->wasNewlyAllocated(HstPtrBegin)) {
// For "PRESENT" entries, we may have cases like the following:
- // map(alloc: p[0]) map(present, alloc: p[0])
- // If the compiler does not merge these entries, then the "PRESENT" entry
- // may be encountered after a previous entry allocated new storage for it.
- // To catch such cases, we should also look at any existing allocations
- // and error out if we have one matching the pointer. We don't need to
- // worry about cases like:
- // map(alloc: p[1:10]) map(present, alloc: p[2:5])
- // as the list-items share storage, but are not identical, which is a
- // user error as per OpenMP.
+ // int *xp = &x[0];
+ // map(alloc: x[:]) map(present, alloc: xp[1])
+ // The "PRESENT" entry may be encountered after a previous entry
+ // allocated new storage for the pointer.
+ // To catch such cases, we need to look at any existing allocations
+ // and error out if we have any matching the pointer.
MESSAGE("device mapping required by 'present' map type modifier does not "
"exist for host address " DPxMOD " (%" PRId64 " bytes)\n",
DPxPTR(HstPtrBegin), DataSize);
@@ -811,14 +808,7 @@ int processAttachEntries(DeviceTy &Device, StateInfoTy &StateInfo,
// Lambda to check if a pointer was newly allocated
auto WasNewlyAllocated = [&](void *Ptr, const char *PtrName) {
- bool IsNewlyAllocated =
- llvm::any_of(StateInfo.NewAllocations, [&](const auto &Alloc) {
- void *AllocPtr = Alloc.first;
- int64_t AllocSize = Alloc.second;
- return Ptr >= AllocPtr &&
- Ptr < reinterpret_cast<void *>(
- reinterpret_cast<char *>(AllocPtr) + AllocSize);
- });
+ bool IsNewlyAllocated = StateInfo.wasNewlyAllocated(Ptr);
DP("Attach %s " DPxMOD " was newly allocated: %s\n", PtrName, DPxPTR(Ptr),
IsNewlyAllocated ? "yes" : "no");
return IsNewlyAllocated;
>From 0e3128f66e59da518fe3ad88fc70b22593dbaacd Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Thu, 13 Nov 2025 12:34:04 -0800
Subject: [PATCH 11/16] Support multiple FROMs falling within the same LAST
entry.
---
offload/libomptarget/OpenMP/Mapping.cpp | 6 +-
offload/libomptarget/omptarget.cpp | 67 ++++++++++++-------
...ring_ptee_tgt_alloc_mapper_alloc_from_to.c | 4 --
...p_ordering_tgt_alloc_from_to_assumedsize.c | 26 +++++++
..._alloc_from_to_structmembers_assumedsize.c | 45 +++++++++++++
...map_ordering_tgt_exit_data_always_always.c | 8 +--
.../map_ordering_tgt_exit_data_delete_from.c | 8 +--
...ng_tgt_exit_data_delete_from_assumedsize.c | 17 +++--
...ng_tgt_exit_data_from_delete_assumedsize.c | 17 +++--
9 files changed, 137 insertions(+), 61 deletions(-)
create mode 100644 offload/test/mapping/map_ordering_tgt_alloc_from_to_assumedsize.c
create mode 100644 offload/test/mapping/map_ordering_tgt_alloc_from_to_structmembers_assumedsize.c
diff --git a/offload/libomptarget/OpenMP/Mapping.cpp b/offload/libomptarget/OpenMP/Mapping.cpp
index 3cfe0c57f0d44..234897421aec3 100644
--- a/offload/libomptarget/OpenMP/Mapping.cpp
+++ b/offload/libomptarget/OpenMP/Mapping.cpp
@@ -325,9 +325,9 @@ TargetPointerResultTy MappingInfoTy::getTargetPointer(
// Lambda to check if this pointer was newly allocated on the current region.
// This is needed to handle cases when the TO entry is encounter after an
- // alloc entry for the same pointer, which increased the ref-count from 0 to
- // 1, has already been encountered before. But because the ref-count was
- // already 1 when TO was encountered, it wouldn't incur a transfer. e.g.
+ // alloc entry for the same pointer. In such cases, the ref-count is already
+ // non-zero when TO is encountered, but we still need to do a transfer. e.g.
+ //
// int *xp = &x[0];
// ... map(alloc: x[:]) map(to: xp[1]).
auto WasNewlyAllocatedForCurrentRegion = [&]() {
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index 0cf85b5d990b8..95e8e05a51693 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -1133,8 +1133,6 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
// Move data back to the host
const bool HasAlways = ArgTypes[I] & OMP_TGT_MAPTYPE_ALWAYS;
const bool HasFrom = ArgTypes[I] & OMP_TGT_MAPTYPE_FROM;
- const bool IsMemberOf = ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF;
- int64_t TransferSize = DataSize; // Size for FROM data-transfer.
// Lambda to perform the actual FROM data retrieval from device to host
auto PerformFromRetrieval = [&](void *HstPtr, void *TgtPtr, int64_t Size,
@@ -1179,6 +1177,12 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
// ... map(delete: p1[:]) map(from: p2[0:1])
// The ref-count becomes zero before encountering the FROM entry, but we
// still need to do a transfer, if it went from non-zero to zero.
+ //
+ // OpenMP 6.0, sec. 7.9.6 "map Clause", p. 284 L24-26:
+ // If the reference count of the corresponding list item is one or if
+ // the always-modifier or delete-modifier is specified, and if the map
+ // type is from, the original list item is updated as if the list item
+ // appeared in a from clause on a target_update directive.
auto WasPreviouslyMarkedForDeletion = [&]() -> bool {
// Check if this pointer falls within the range of any deleted entry
for (const auto &DeleteEntry : StateInfo->DeleteEntries) {
@@ -1215,61 +1219,72 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
// Then we can forego tacking TransferredFromPtrs.
StateInfo->TransferredFromPtrs.insert(HstPtrBegin);
- Ret = PerformFromRetrieval(HstPtrBegin, TgtPtrBegin, TransferSize,
+ Ret = PerformFromRetrieval(HstPtrBegin, TgtPtrBegin, DataSize,
TPR.getEntry());
if (Ret != OFFLOAD_SUCCESS)
return OFFLOAD_FAIL;
} else if (!FromCopyBackAlreadyDone && IsMapFromOnNonHostNonZeroData &&
- !IsLastOrHasAlwaysOrWasForceDeleted() && !IsMemberOf) {
+ !IsLastOrHasAlwaysOrWasForceDeleted()) {
// We can have cases like the following:
- // ... map(storage: p1[0:1]) map(from: p1[0:1])
+ // int *xp = &x[0];
+ // ... map(storage: x[:]) map(from: xp[1:1])
//
// where it's possible that when the FROM entry is processed, the
// ref count is not zero, so no data transfer happens for it. But
- // the ref-count can go down to zero by the end of the directive
+ // the ref-count can go down to zero once all maps have been processed
// in which case a transfer should happen.
//
// So, we keep track of any skipped FROM data-transfers, in case
// the ref-count goes down to zero later on.
//
- // This should be limited to non-member-of entries because for member-of,
- // their ref-count should go down only once as part of the parent.
+ // This cannot be handled in the compiler for all cases because the
+ // list-items may look very different, as shown in the example above,
+ // which is allowed with OpenMP 6.0:
+ //
+ // OpenMP 6.0, sec. 7.9.6 "map Clause", p. 286 L18-21:
+ // Two list items of the map clauses on the same construct must not share
+ // original storage unless one of the following is true: they are the same
+ // list item, one is the containing structure of the other, at least one
+ // is an assumed-size array, or at least one is implicitly mapped due to
+ // the list item also appearing in a use_device_addr clause.
StateInfo->SkippedFromEntries[HstPtrBegin] = DataSize;
DP("Skipping FROM map transfer for HstPtr=" DPxMOD ", Size=%" PRId64 "\n",
DPxPTR(HstPtrBegin), DataSize);
} else if (!FromCopyBackAlreadyDone && TPR.Flags.IsLast) {
- // Even if this is not a FROM or DELETE entry, if the ref-count went to
+ // Even if this is not a FROM entry, if the ref-count went to
// zero (IsLast=true), we should perform any previously skipped FROM
// transfers that fall within this entry's range.
SmallVector<void *, 32> ToRemove;
- for (auto &SkippedEntry : StateInfo->SkippedFromEntries) {
- void *SkippedPtr = SkippedEntry.first;
- int64_t SkippedSize = SkippedEntry.second;
- uintptr_t SkippedBegin = (uintptr_t)SkippedPtr;
+ for (auto &SkippedFromEntry : StateInfo->SkippedFromEntries) {
+ void *FromBeginPtr = SkippedFromEntry.first;
+ int64_t FromDataSize = SkippedFromEntry.second;
+ uintptr_t FromBeginPtrInt = (uintptr_t)FromBeginPtr;
- uintptr_t EntryBegin = TPR.getEntry()->HstPtrBegin;
- uintptr_t EntryEnd = TPR.getEntry()->HstPtrEnd;
+ uintptr_t DeleteBeginPtrInt = TPR.getEntry()->HstPtrBegin;
+ uintptr_t DeleteEndPtrInt = TPR.getEntry()->HstPtrEnd;
// Check if skipped entry overlaps or is contained within current entry
- if (SkippedBegin >= EntryBegin && SkippedBegin < EntryEnd) {
- DP("Found skipped FROM within deleted region: HstPtr=" DPxMOD
- " size=%" PRId64 " within [" DPxMOD ", " DPxMOD ")\n",
- DPxPTR(SkippedPtr), SkippedSize, DPxPTR(EntryBegin),
- DPxPTR(EntryEnd));
+ if (FromBeginPtrInt >= DeleteBeginPtrInt &&
+ FromBeginPtrInt < DeleteEndPtrInt) {
+ DP("Found skipped FROM entry: HstPtr=" DPxMOD " size=%" PRId64
+ " within region being deleted [" DPxMOD ", " DPxMOD ")\n",
+ DPxPTR(FromBeginPtr), FromDataSize, DPxPTR(DeleteBeginPtrInt),
+ DPxPTR(DeleteEndPtrInt));
// Calculate offset within the target pointer
- int64_t Offset = SkippedBegin - EntryBegin;
- void *SkippedTgtPtr = (void *)((char *)TgtPtrBegin + Offset);
+ int64_t Offset = FromBeginPtrInt - DeleteBeginPtrInt;
+ void *FromTgtBeginPtr = (void *)((char *)TgtPtrBegin + Offset);
// Perform the retrieval for this skipped entry
- int Ret = PerformFromRetrieval((void *)SkippedBegin, SkippedTgtPtr,
- SkippedSize, TPR.getEntry());
+ int Ret =
+ PerformFromRetrieval((void *)FromBeginPtrInt, FromTgtBeginPtr,
+ FromDataSize, TPR.getEntry());
if (Ret != OFFLOAD_SUCCESS)
return OFFLOAD_FAIL;
- StateInfo->TransferredFromPtrs.insert(SkippedPtr);
- ToRemove.push_back(SkippedPtr);
+ StateInfo->TransferredFromPtrs.insert(FromBeginPtr);
+ ToRemove.push_back(FromBeginPtr);
}
}
diff --git a/offload/test/mapping/map_ordering_ptee_tgt_alloc_mapper_alloc_from_to.c b/offload/test/mapping/map_ordering_ptee_tgt_alloc_mapper_alloc_from_to.c
index 3ad8748cb0572..1ddc241b429a1 100644
--- a/offload/test/mapping/map_ordering_ptee_tgt_alloc_mapper_alloc_from_to.c
+++ b/offload/test/mapping/map_ordering_ptee_tgt_alloc_mapper_alloc_from_to.c
@@ -7,10 +7,6 @@
// construct, even if the ref-count of the pointee x has not gone down to 0
// when "from" is encountered.
-// This currently fails, but should start passing once ATTACH-style maps are
-// enabled for mappers (#166874).
-// XFAIL: *
-
#include <stdio.h>
typedef struct {
diff --git a/offload/test/mapping/map_ordering_tgt_alloc_from_to_assumedsize.c b/offload/test/mapping/map_ordering_tgt_alloc_from_to_assumedsize.c
new file mode 100644
index 0000000000000..a272732e59343
--- /dev/null
+++ b/offload/test/mapping/map_ordering_tgt_alloc_from_to_assumedsize.c
@@ -0,0 +1,26 @@
+// RUN: %libomptarget-compile-generic -fopenmp-version=60
+
+// The test ensures that even though the from/to maps are on
+// a list-item that looks different from the "alloc" list-items
+// on assumed-size arrays, that are encountered before the to/from
+// entries at run time, the transfers still kick in.
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+ int x[10];
+ x[1] = 111;
+
+ int *xp1, *xp2;
+ xp1 = xp2 = &x[0];
+
+#pragma omp target map(alloc : x[ : ]) map(from : xp1[1]) map(to : xp1[1]) \
+ map(alloc : xp2[ : ])
+ {
+ printf("%d\n", x[1]); // CHECK: 111
+ x[1] = x[1] + 111;
+ }
+
+ printf("%d\n", x[1]); // CHECK: 222
+}
diff --git a/offload/test/mapping/map_ordering_tgt_alloc_from_to_structmembers_assumedsize.c b/offload/test/mapping/map_ordering_tgt_alloc_from_to_structmembers_assumedsize.c
new file mode 100644
index 0000000000000..822c68cc673fe
--- /dev/null
+++ b/offload/test/mapping/map_ordering_tgt_alloc_from_to_structmembers_assumedsize.c
@@ -0,0 +1,45 @@
+// RUN: %libomptarget-compile-generic -fopenmp-version=60
+// RUN: env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefix=DEBUG -check-prefix=CHECK
+// REQUIRES: libomptarget-debug
+
+#include <omp.h>
+#include <stdio.h>
+
+// This test shows that when an "release" is being handled, we need to look at
+// all previously skipped data-transfers within the range of the memory being
+// released. Here, when "s[:]" is being released, we need to honor the two
+// "from" on sp1->x and sp1->y.
+
+typedef struct {
+ int x;
+ int y;
+} S;
+
+int main() {
+ S s[10];
+ s[1].x = 111;
+ s[1].y = 111;
+
+ S *sp1 = &s[1];
+
+#pragma omp target map(alloc : s[ : ]) map(from : sp1 -> x, sp1->y) \
+ map(to : sp1->x, sp1->y) map(alloc : s[ : ])
+ {
+ fprintf(stderr, "%d %d\n", s[1].x, s[1].y); // CHECK: 111 111
+ s[1].x = s[1].x + 111;
+ s[1].y = s[1].y + 111;
+ }
+
+ // DEBUG: omptarget --> Found skipped FROM entry:
+ // DEBUG-SAME: HstPtr=0x[[#%x,HOST_ADDRX:]] size=[[#%u,SIZE:]]
+ // DEBUG-SAME: within region being deleted
+ // DEBUG: omptarget --> Moving [[#SIZE]] bytes (tgt:0x{{.*}}) ->
+ // DEBUG-SAME: (hst:0x{{0*}}[[#HOST_ADDRX]])
+ // DEBUG: omptarget --> Found skipped FROM entry:
+ // DEBUG-SAME: HstPtr=0x[[#%x,HOST_ADDRY:]] size=[[#%u,SIZE:]]
+ // DEBUG-SAME: within region being deleted
+ // DEBUG: omptarget --> Moving [[#SIZE]] bytes (tgt:0x{{.*}}) ->
+ // DEBUG-SAME: (hst:0x{{0*}}[[#HOST_ADDRY]])
+ fprintf(stderr, "%d %d\n", s[1].x, s[1].y); // CHECK: 222 222
+}
diff --git a/offload/test/mapping/map_ordering_tgt_exit_data_always_always.c b/offload/test/mapping/map_ordering_tgt_exit_data_always_always.c
index 43f3aae26b009..0e00f9efc8289 100644
--- a/offload/test/mapping/map_ordering_tgt_exit_data_always_always.c
+++ b/offload/test/mapping/map_ordering_tgt_exit_data_always_always.c
@@ -13,12 +13,10 @@ int main() {
#pragma omp target data map(alloc : x)
{
#pragma omp target enter data map(alloc : x) map(to : x)
- {
#pragma omp target map(present, alloc : x)
- {
- printf("In tgt: %d\n", x); // CHECK-NOT: In tgt: 111
- x = 222;
- }
+ {
+ printf("In tgt: %d\n", x); // CHECK-NOT: In tgt: 111
+ x = 222;
}
#pragma omp target exit data map(always, from : x) map(always, from : x)
// DEBUG: omptarget --> Moving 4 bytes (tgt:0x{{.*}}) -> (hst:0x{{.*}})
diff --git a/offload/test/mapping/map_ordering_tgt_exit_data_delete_from.c b/offload/test/mapping/map_ordering_tgt_exit_data_delete_from.c
index ad7db66890edb..896bc8bb18efa 100644
--- a/offload/test/mapping/map_ordering_tgt_exit_data_delete_from.c
+++ b/offload/test/mapping/map_ordering_tgt_exit_data_delete_from.c
@@ -7,12 +7,10 @@ int main() {
#pragma omp target data map(alloc : x)
{
#pragma omp target enter data map(alloc : x) map(to : x)
- {
#pragma omp target map(present, alloc : x)
- {
- printf("%d\n", x); // CHECK-NOT: 111
- x = 222;
- }
+ {
+ printf("%d\n", x); // CHECK-NOT: 111
+ x = 222;
}
#pragma omp target exit data map(delete : x) map(from : x) map(delete : x)
printf("%d\n", x); // CHECK: 222
diff --git a/offload/test/mapping/map_ordering_tgt_exit_data_delete_from_assumedsize.c b/offload/test/mapping/map_ordering_tgt_exit_data_delete_from_assumedsize.c
index f8a117fc0ecf7..1e4cef4b66335 100644
--- a/offload/test/mapping/map_ordering_tgt_exit_data_delete_from_assumedsize.c
+++ b/offload/test/mapping/map_ordering_tgt_exit_data_delete_from_assumedsize.c
@@ -20,18 +20,17 @@ int main() {
#pragma omp target data map(alloc : x)
{
#pragma omp target enter data map(alloc : x) map(to : x)
- {
#pragma omp target map(present, alloc : x)
- {
- printf("In tgt: %d\n", x[0]); // CHECK-NOT: In tgt: 111
- x[0] = 222;
- }
+ {
+ printf("In tgt: %d\n", x[1]); // CHECK-NOT: In tgt: 111
+ x[1] = 222;
}
-// DEBUG: omptarget --> Found previously skipped FROM transfer
-// DEBUG-SAME: for HstPtr=0x[[#%x,HOST_ADDR:]], with size [[#%u,SIZE:]]
+// DEBUG: omptarget --> Found skipped FROM entry
+// DEBUG-SAME: HstPtr=0x[[#%x,HOST_ADDR:]] size=[[#%u,SIZE:]]
+// DEBUG-SAME: within region being deleted
// DEBUG: omptarget --> Moving [[#SIZE]] bytes
// DEBUG-SAME: (tgt:0x{{.*}}) -> (hst:0x{{0*}}[[#HOST_ADDR]])
-#pragma omp target exit data map(delete : p1x[ : ]) map(from : p2x[0])
- printf("%d\n", x[0]); // CHECK: 222
+#pragma omp target exit data map(delete : p1x[ : ]) map(from : p2x[1])
+ printf("%d\n", x[1]); // CHECK: 222
}
}
diff --git a/offload/test/mapping/map_ordering_tgt_exit_data_from_delete_assumedsize.c b/offload/test/mapping/map_ordering_tgt_exit_data_from_delete_assumedsize.c
index 3f5b08d3473f8..d939bb66235c1 100644
--- a/offload/test/mapping/map_ordering_tgt_exit_data_from_delete_assumedsize.c
+++ b/offload/test/mapping/map_ordering_tgt_exit_data_from_delete_assumedsize.c
@@ -16,23 +16,22 @@
int main() {
int x[10];
int *p1x, *p2x;
- p1x = p2x = &x[0];
+ p1x = p2x = &x[1];
+ x[1] = 111;
#pragma omp target data map(alloc : x)
{
#pragma omp target enter data map(alloc : x) map(to : x)
- {
#pragma omp target map(present, alloc : x)
- {
- printf("In tgt: %d\n", x[0]); // CHECK-NOT: In tgt: 111
- x[0] = 222;
- }
+ {
+ fprintf(stderr, "In tgt: %d\n", x[1]); // CHECK-NOT: In tgt: 111
+ x[1] = 222;
}
- // DEBUG: omptarget --> Pointer HstPtr=0x[[#%x,HOST_ADDR:]]
- // DEBUG-SAME: was previously marked for deletion
+ // DEBUG: omptarget --> Pointer HstPtr=0x[[#%x,HOST_ADDR:]] falls within a
+ // DEBUG-SAME: range previously marked for deletion
// DEBUG: omptarget --> Moving {{.*}} bytes
// DEBUG-SAME: (tgt:0x{{.*}}) -> (hst:0x{{0*}}[[#HOST_ADDR]])
#pragma omp target exit data map(from : p2x[0]) map(delete : p1x[ : ])
- printf("%d\n", x[0]); // CHECK: 222
+ fprintf(stderr, "%d\n", x[1]); // CHECK: 222
}
}
>From a49f89dbbbd5320963d824b045fd638eef069813 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Wed, 4 Feb 2026 15:52:23 -0800
Subject: [PATCH 12/16] Address Robert's review comments.
---
offload/include/OpenMP/Mapping.h | 4 ++--
offload/libomptarget/omptarget.cpp | 2 +-
...map_ordering_tgt_alloc_from_to_structmembers_assumedsize.c | 4 +++-
3 files changed, 6 insertions(+), 4 deletions(-)
diff --git a/offload/include/OpenMP/Mapping.h b/offload/include/OpenMP/Mapping.h
index 7840c8f2dd598..e17a13f4917f1 100644
--- a/offload/include/OpenMP/Mapping.h
+++ b/offload/include/OpenMP/Mapping.h
@@ -533,8 +533,8 @@ struct StateInfoTy {
void *AllocPtr = Alloc.first;
int64_t AllocSize = Alloc.second;
return Ptr >= AllocPtr &&
- Ptr < reinterpret_cast<void *>(
- reinterpret_cast<char *>(AllocPtr) + AllocSize);
+ Ptr < static_cast<void *>(static_cast<char *>(AllocPtr) +
+ AllocSize);
});
}
};
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index 11efa82a1c3fc..586d0fb7178ae 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -517,7 +517,7 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
"handling ATTACH and TO/TOFROM map-types.");
// process each input.
for (int32_t I = 0; I < ArgNum; ++I) {
- // Ignore private variables and arrays - there is no mapping for t.attahem.
+ // Ignore private variables and arrays - there is no mapping for them.
if ((ArgTypes[I] & OMP_TGT_MAPTYPE_LITERAL) ||
(ArgTypes[I] & OMP_TGT_MAPTYPE_PRIVATE))
continue;
diff --git a/offload/test/mapping/map_ordering_tgt_alloc_from_to_structmembers_assumedsize.c b/offload/test/mapping/map_ordering_tgt_alloc_from_to_structmembers_assumedsize.c
index 822c68cc673fe..4be33a57eeba4 100644
--- a/offload/test/mapping/map_ordering_tgt_alloc_from_to_structmembers_assumedsize.c
+++ b/offload/test/mapping/map_ordering_tgt_alloc_from_to_structmembers_assumedsize.c
@@ -1,6 +1,8 @@
// RUN: %libomptarget-compile-generic -fopenmp-version=60
+// RUN: %libomptarget-run-generic
+// RUN: | %fcheck-generic -check-prefix=CHECK
// RUN: env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 \
-// RUN: | %fcheck-generic -check-prefix=DEBUG -check-prefix=CHECK
+// RUN: | %fcheck-generic -check-prefix=DEBUG
// REQUIRES: libomptarget-debug
#include <omp.h>
>From 50f040097bfd6b70dddde1937fc4e131d16918d7 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Thu, 5 Feb 2026 14:53:51 -0800
Subject: [PATCH 13/16] Clean-up some comments, tests.
---
offload/include/OpenMP/Mapping.h | 6 ++--
offload/libomptarget/OpenMP/Mapping.cpp | 12 ++++++--
offload/libomptarget/omptarget.cpp | 23 +++++++-------
...ring_ptee_tgt_alloc_mapper_alloc_from_to.c | 21 ++++++++++---
..._alloc_tgt_mapper_present_delete_from_to.c | 20 +++++++++----
.../mapping/map_ordering_tgt_alloc_from_to.c | 8 ++---
...p_ordering_tgt_alloc_from_to_assumedsize.c | 10 +++++--
..._alloc_from_to_structmembers_assumedsize.c | 30 ++++++++++---------
.../map_ordering_tgt_alloc_present_tofrom.c | 13 ++++----
...map_ordering_tgt_exit_data_always_always.c | 4 +--
.../map_ordering_tgt_exit_data_delete_from.c | 1 +
...ng_tgt_exit_data_delete_from_assumedsize.c | 15 ++++++----
...ng_tgt_exit_data_from_delete_assumedsize.c | 20 +++++++------
13 files changed, 110 insertions(+), 73 deletions(-)
diff --git a/offload/include/OpenMP/Mapping.h b/offload/include/OpenMP/Mapping.h
index e17a13f4917f1..b052504b425e3 100644
--- a/offload/include/OpenMP/Mapping.h
+++ b/offload/include/OpenMP/Mapping.h
@@ -511,7 +511,7 @@ struct StateInfoTy {
/// Key: host pointer, Value: data size.
llvm::DenseMap<void *, int64_t> SkippedFromEntries;
- /// Host pointers for which we have attempted a FROM transfer at some point
+ /// Host pointers for which we have triggered a FROM transfer at some point
/// during targetDataEnd. Used to avoid duplicate transfers.
llvm::SmallSet<void *, 32> TransferredFromPtrs;
@@ -528,8 +528,8 @@ struct StateInfoTy {
/// Check if a pointer falls within any of the newly allocated ranges.
/// Returns true if the pointer is within a newly allocated region.
bool wasNewlyAllocated(void *Ptr) const {
- return std::any_of(
- NewAllocations.begin(), NewAllocations.end(), [&](const auto &Alloc) {
+ return llvm::any_of(
+ NewAllocations, [&](const auto &Alloc) {
void *AllocPtr = Alloc.first;
int64_t AllocSize = Alloc.second;
return Ptr >= AllocPtr &&
diff --git a/offload/libomptarget/OpenMP/Mapping.cpp b/offload/libomptarget/OpenMP/Mapping.cpp
index d8725e01cc50b..6f075f6545b3d 100644
--- a/offload/libomptarget/OpenMP/Mapping.cpp
+++ b/offload/libomptarget/OpenMP/Mapping.cpp
@@ -330,12 +330,18 @@ TargetPointerResultTy MappingInfoTy::getTargetPointer(
HDTTMap.destroy();
// Lambda to check if this pointer was newly allocated on the current region.
- // This is needed to handle cases when the TO entry is encounter after an
+ // This is needed to handle cases when the TO entry is encountered after an
// alloc entry for the same pointer. In such cases, the ref-count is already
// non-zero when TO is encountered, but we still need to do a transfer. e.g.
//
- // int *xp = &x[0];
- // ... map(alloc: x[:]) map(to: xp[1]).
+ // struct S {
+ // int *p;
+ // };
+ // #pragma omp declare mapper(id : S s) map(to: s.p, s.p[0 : 10])
+ //
+ // S s1;
+ // ...
+ // #pragma omp target map(alloc : s1.p[0 : 10]) map(mapper(id), to : s1)
auto WasNewlyAllocatedForCurrentRegion = [&]() {
if (!StateInfo)
return false;
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index 586d0fb7178ae..b5b4cf2d984b0 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -1131,7 +1131,6 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
HstPtrBegin, DataSize, UpdateRef, HasHoldModifier, !IsImplicit,
ForceDelete, /*FromDataEnd=*/true);
void *TgtPtrBegin = TPR.TargetPointer;
-
if (!TPR.isPresent() && !TPR.isHostPointer() &&
(DataSize || HasPresentModifier)) {
ODBG(ODT_Mapping) << "Mapping does not exist ("
@@ -1247,11 +1246,11 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
void *DeletePtr = DeleteEntry.first;
int64_t DeleteSize = DeleteEntry.second;
if (HstPtrBegin >= DeletePtr &&
- HstPtrBegin < (void *)((char *)DeletePtr + DeleteSize)) {
+ HstPtrBegin < static_cast<void *>(static_cast<char *>(DeletePtr) + DeleteSize)) {
ODBG(ODT_Mapping)
<< "Pointer HstPtr=" << HstPtrBegin
<< " falls within a range previously marked for deletion ["
- << DeletePtr << ", " << (char *)DeletePtr + DeleteSize
+ << DeletePtr << ", " << static_cast<char *>(DeletePtr) + DeleteSize
<< ") with size=" << DeleteSize;
return true;
}
@@ -1281,16 +1280,16 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
TPR.getEntry());
if (Ret != OFFLOAD_SUCCESS)
return OFFLOAD_FAIL;
- } else if (!FromCopyBackAlreadyDone && IsMapFromOnNonHostNonZeroData &&
- !IsLastOrHasAlwaysOrWasForceDeleted()) {
+ } else if (!FromCopyBackAlreadyDone && (IsMapFromOnNonHostNonZeroData &&
+ !IsLastOrHasAlwaysOrWasForceDeleted())) {
// We can have cases like the following:
- // int *xp = &x[0];
- // ... map(storage: x[:]) map(from: xp[1:1])
+ // p1 = p2 = &x;
+ // ... map(storage: p1[:]) map(from: p2[1:1])
//
// where it's possible that when the FROM entry is processed, the
// ref count is not zero, so no data transfer happens for it. But
// the ref-count can go down to zero once all maps have been processed
- // in which case a transfer should happen.
+ // for the current construct, in which case a transfer should happen.
//
// So, we keep track of any skipped FROM data-transfers, in case
// the ref-count goes down to zero later on.
@@ -1317,7 +1316,7 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
for (auto &SkippedFromEntry : StateInfo->SkippedFromEntries) {
void *FromBeginPtr = SkippedFromEntry.first;
int64_t FromDataSize = SkippedFromEntry.second;
- uintptr_t FromBeginPtrInt = (uintptr_t)FromBeginPtr;
+ uintptr_t FromBeginPtrInt = reinterpret_cast<uintptr_t>(FromBeginPtr);
uintptr_t DeleteBeginPtrInt = TPR.getEntry()->HstPtrBegin;
uintptr_t DeleteEndPtrInt = TPR.getEntry()->HstPtrEnd;
@@ -1328,15 +1327,15 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
ODBG(ODT_Mapping)
<< "Found skipped FROM entry: HstPtr=" << FromBeginPtr
<< " size=" << FromDataSize << " within region being deleted ["
- << DeleteBeginPtrInt << ", " << DeleteEndPtrInt << ")";
+ << reinterpret_cast<void*>(DeleteBeginPtrInt) << ", " << reinterpret_cast<void*>(DeleteEndPtrInt) << ")";
// Calculate offset within the target pointer
int64_t Offset = FromBeginPtrInt - DeleteBeginPtrInt;
- void *FromTgtBeginPtr = (void *)((char *)TgtPtrBegin + Offset);
+ void *FromTgtBeginPtr = static_cast<void *>(static_cast<char *>(TgtPtrBegin) + Offset);
// Perform the retrieval for this skipped entry
int Ret =
- PerformFromRetrieval((void *)FromBeginPtrInt, FromTgtBeginPtr,
+ PerformFromRetrieval(reinterpret_cast<void *>(FromBeginPtrInt), FromTgtBeginPtr,
FromDataSize, TPR.getEntry());
if (Ret != OFFLOAD_SUCCESS)
return OFFLOAD_FAIL;
diff --git a/offload/test/mapping/map_ordering_ptee_tgt_alloc_mapper_alloc_from_to.c b/offload/test/mapping/map_ordering_ptee_tgt_alloc_mapper_alloc_from_to.c
index 1ddc241b429a1..9a94cab804b71 100644
--- a/offload/test/mapping/map_ordering_ptee_tgt_alloc_mapper_alloc_from_to.c
+++ b/offload/test/mapping/map_ordering_ptee_tgt_alloc_mapper_alloc_from_to.c
@@ -1,7 +1,12 @@
-// RUN: %libomptarget-compile-run-and-check-generic
+// RUN: %libomptarget-compile-generic
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefix=CHECK
+// RUN: env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefix=DEBUG
+// REQUIRES: libomptarget-debug
-// Since the allocation of the pointee happens as on the "target" construct (1),
-// the "to" transfer requested as part of the mapper should also happen.
+// Since the allocation of the pointee happens on the "target" construct (1),
+// the "to" transfer requested as part of the mapper (2) should also happen.
//
// Similarly, the "from" transfer should also happen at the end of the target
// construct, even if the ref-count of the pointee x has not gone down to 0
@@ -14,7 +19,7 @@ typedef struct {
int *q;
} S;
#pragma omp declare mapper(my_mapper : S s) map(alloc : s.p, s.p[0 : 10]) \
- map(from : s.p[0 : 10]) map(to : s.p[0 : 10]) map(alloc : s.p[0 : 10])
+ map(from : s.p[0 : 10]) map(to : s.p[0 : 10]) map(alloc : s.p[0 : 10]) // (2)
S s1;
int main() {
@@ -22,6 +27,10 @@ int main() {
x[1] = 111;
s1.q = s1.p = &x[0];
+// clang-format off
+// DEBUG: omptarget --> HstPtrBegin 0x[[#%x,HOST_ADDRX:]] was newly allocated for the current region
+// DEBUG: omptarget --> Moving [[#%u,SIZEX:]] bytes (hst:0x{{0*}}[[#HOST_ADDRX]]) -> (tgt:0x{{.*}})
+// clang-format on
#pragma omp target map(alloc : s1.p[0 : 10]) \
map(mapper(my_mapper), tofrom : s1) // (1)
{
@@ -29,5 +38,9 @@ int main() {
s1.p[1] = s1.p[1] + 111;
}
+// clang-format off
+// DEBUG: omptarget --> Found skipped FROM entry: HstPtr=0x{{0*}}[[#HOST_ADDRX]] size=[[#SIZEX]] within region being deleted
+// DEBUG: omptarget --> Moving [[#SIZEX]] bytes (tgt:0x{{.*}}) -> (hst:0x{{0*}}[[#HOST_ADDRX]])
+// clang-format on
printf("%d\n", s1.p[1]); // CHECK: 222
}
diff --git a/offload/test/mapping/map_ordering_ptee_tgt_data_alloc_tgt_mapper_present_delete_from_to.c b/offload/test/mapping/map_ordering_ptee_tgt_data_alloc_tgt_mapper_present_delete_from_to.c
index 5fb196e31d9c2..791ff3ddb2a3b 100644
--- a/offload/test/mapping/map_ordering_ptee_tgt_data_alloc_tgt_mapper_present_delete_from_to.c
+++ b/offload/test/mapping/map_ordering_ptee_tgt_data_alloc_tgt_mapper_present_delete_from_to.c
@@ -1,15 +1,19 @@
-// RUN: %libomptarget-compile-generic -fopenmp-version=60
-// RUN: %libomptarget-run-generic | %fcheck-generic
+// RUN: %libomptarget-compile-generic
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefix=CHECK
+// RUN: env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefix=DEBUG
+// REQUIRES: libomptarget-debug
-// The "present" check should pass on the "target" construct // (2),
+// The "present" check should pass on the "target" construct (2),
// and there should be no "to" transfer, because the pointee "x" is already
// present (because of (1)).
// However, there should be a "from" transfer at the end of (2) because of the
// "delete" on the mapper.
-// This currently fails, but should start passing once ATTACH-style maps are
+// FIXME: This currently fails, but should start passing once ATTACH-style maps are
// enabled for mappers (#166874).
-// XFAIL: *
+// UNSUPPORTED: true
#include <stdio.h>
@@ -29,11 +33,17 @@ int main() {
#pragma omp target data map(alloc : x) // (1)
{
+// DEBUG-NOT: omptarget --> Moving {{.*}} bytes (hst:0x{{.*}}) -> (tgt:0x{{.*}})
#pragma omp target map(mapper(my_mapper), tofrom : s1) // (2)
{
+ // NOTE: It's ok for this to be 111 under "unified_shared_memory"
printf("%d\n", s1.p[1]); // CHECK-NOT: 111
s1.p[1] = 222;
}
printf("%d\n", s1.p[1]); // CHECK: 222
}
+// clang-format off
+// DEBUG: omptarget --> Found skipped FROM entry: HstPtr=0x[[#%x,HOST_ADDR:]] size=[[#%u,SIZE:]] within region being deleted
+// DEBUG: omptarget --> Moving [[#SIZE]] bytes (tgt:0x{{.*}}) -> (hst:0x{{0*}}[[#HOST_ADDR]])
+// clang-format on
}
diff --git a/offload/test/mapping/map_ordering_tgt_alloc_from_to.c b/offload/test/mapping/map_ordering_tgt_alloc_from_to.c
index 71f68134c7317..25f24b37c2bc2 100644
--- a/offload/test/mapping/map_ordering_tgt_alloc_from_to.c
+++ b/offload/test/mapping/map_ordering_tgt_alloc_from_to.c
@@ -11,10 +11,10 @@
int main() {
int x = 111;
- // DEBUG: omptarget --> HstPtrBegin 0x[[#%x,HOST_ADDR:]] was newly allocated
- // DEBUG-SAME: for the current region
- // DEBUG: omptarget --> Moving {{.*}} bytes
- // DEBUG-SAME: (hst:0x{{0*}}[[#HOST_ADDR]]) -> (tgt:0x{{.*}})
+ // clang-format off
+ // DEBUG: omptarget --> HstPtrBegin 0x[[#%x,HOST_ADDR:]] was newly allocated for the current region
+ // DEBUG: omptarget --> Moving {{.*}} bytes (hst:0x{{0*}}[[#HOST_ADDR]]) -> (tgt:0x{{.*}})
+ // clang-format on
#pragma omp target map(alloc : x) map(from : x) map(to : x) map(alloc : x)
{
printf("%d\n", x); // CHECK: 111
diff --git a/offload/test/mapping/map_ordering_tgt_alloc_from_to_assumedsize.c b/offload/test/mapping/map_ordering_tgt_alloc_from_to_assumedsize.c
index a272732e59343..5612aa5cb01c1 100644
--- a/offload/test/mapping/map_ordering_tgt_alloc_from_to_assumedsize.c
+++ b/offload/test/mapping/map_ordering_tgt_alloc_from_to_assumedsize.c
@@ -1,9 +1,13 @@
// RUN: %libomptarget-compile-generic -fopenmp-version=60
+// RUN: %libomptarget-run-generic | %fcheck-generic
// The test ensures that even though the from/to maps are on
-// a list-item that looks different from the "alloc" list-items
-// on assumed-size arrays, that are encountered before the to/from
-// entries at run time, the transfers still kick in.
+// a list-item that looks different from the "alloc" list-items,
+// which are encountered before the to/from entries at run time,
+// the transfers still happen.
+//
+// NOTE: This is not fully OpenMP compliant, but is a simpler
+// proxy for when this happens via mappers.
#include <omp.h>
#include <stdio.h>
diff --git a/offload/test/mapping/map_ordering_tgt_alloc_from_to_structmembers_assumedsize.c b/offload/test/mapping/map_ordering_tgt_alloc_from_to_structmembers_assumedsize.c
index 4be33a57eeba4..54fa075dc4827 100644
--- a/offload/test/mapping/map_ordering_tgt_alloc_from_to_structmembers_assumedsize.c
+++ b/offload/test/mapping/map_ordering_tgt_alloc_from_to_structmembers_assumedsize.c
@@ -1,5 +1,5 @@
// RUN: %libomptarget-compile-generic -fopenmp-version=60
-// RUN: %libomptarget-run-generic
+// RUN: %libomptarget-run-generic 2>&1 \
// RUN: | %fcheck-generic -check-prefix=CHECK
// RUN: env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 \
// RUN: | %fcheck-generic -check-prefix=DEBUG
@@ -8,10 +8,10 @@
#include <omp.h>
#include <stdio.h>
-// This test shows that when an "release" is being handled, we need to look at
+// This test shows that when a "release" is being handled, we look at
// all previously skipped data-transfers within the range of the memory being
-// released. Here, when "s[:]" is being released, we need to honor the two
-// "from" on sp1->x and sp1->y.
+// released. Here, when "s[:]" is being released, we honor the two "from"s
+// on sp1->x and sp1->y.
typedef struct {
int x;
@@ -25,6 +25,12 @@ int main() {
S *sp1 = &s[1];
+// clang-format off
+// DEBUG: omptarget --> HstPtrBegin 0x[[#%x,HOST_ADDRX:]] was newly allocated for the current region
+// DEBUG: omptarget --> Moving [[#%u,SIZEX:]] bytes (hst:0x{{0*}}[[#HOST_ADDRX]]) -> (tgt:0x{{.*}})
+// DEBUG: omptarget --> HstPtrBegin 0x[[#%x,HOST_ADDRY:]] was newly allocated for the current region
+// DEBUG: omptarget --> Moving [[#%u,SIZEY:]] bytes (hst:0x{{0*}}[[#HOST_ADDRY]]) -> (tgt:0x{{.*}})
+// clang-format on
#pragma omp target map(alloc : s[ : ]) map(from : sp1 -> x, sp1->y) \
map(to : sp1->x, sp1->y) map(alloc : s[ : ])
{
@@ -33,15 +39,11 @@ int main() {
s[1].y = s[1].y + 111;
}
- // DEBUG: omptarget --> Found skipped FROM entry:
- // DEBUG-SAME: HstPtr=0x[[#%x,HOST_ADDRX:]] size=[[#%u,SIZE:]]
- // DEBUG-SAME: within region being deleted
- // DEBUG: omptarget --> Moving [[#SIZE]] bytes (tgt:0x{{.*}}) ->
- // DEBUG-SAME: (hst:0x{{0*}}[[#HOST_ADDRX]])
- // DEBUG: omptarget --> Found skipped FROM entry:
- // DEBUG-SAME: HstPtr=0x[[#%x,HOST_ADDRY:]] size=[[#%u,SIZE:]]
- // DEBUG-SAME: within region being deleted
- // DEBUG: omptarget --> Moving [[#SIZE]] bytes (tgt:0x{{.*}}) ->
- // DEBUG-SAME: (hst:0x{{0*}}[[#HOST_ADDRY]])
+ // clang-format off
+ // DEBUG: omptarget --> Found skipped FROM entry: HstPtr=0x{{0*}}[[#HOST_ADDRY]] size=[[#SIZEY]] within region being deleted
+ // DEBUG: omptarget --> Moving [[#SIZEY]] bytes (tgt:0x{{.*}}) -> (hst:0x{{0*}}[[#HOST_ADDRY]])
+ // DEBUG: omptarget --> Found skipped FROM entry: HstPtr=0x{{0*}}[[#HOST_ADDRX]] size=[[#SIZEX]] within region being deleted
+ // DEBUG: omptarget --> Moving [[#SIZEX]] bytes (tgt:0x{{.*}}) -> (hst:0x{{0*}}[[#HOST_ADDRX]])
+ // clang-format on
fprintf(stderr, "%d %d\n", s[1].x, s[1].y); // CHECK: 222 222
}
diff --git a/offload/test/mapping/map_ordering_tgt_alloc_present_tofrom.c b/offload/test/mapping/map_ordering_tgt_alloc_present_tofrom.c
index f8f397efc4acd..32d5cf4b1b06b 100644
--- a/offload/test/mapping/map_ordering_tgt_alloc_present_tofrom.c
+++ b/offload/test/mapping/map_ordering_tgt_alloc_present_tofrom.c
@@ -8,16 +8,13 @@ int main() {
// CHECK: addr=0x[[#%x,HOST_ADDR:]], size=[[#%u,SIZE:]]
int x = 111;
fprintf(stderr, "addr=%p, size=%ld\n", &x, sizeof(x));
-// CHECK: omptarget message: device mapping required by 'present' map type
-// modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]]
-// bytes)
-// CHECK: omptarget error: Pointer 0x{{0*}}[[#HOST_ADDR]] was not present
-// on the device upon entry to the region.
-// ('present' map type modifier).
+// clang-format off
+// CHECK: omptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes)
+// CHECK: omptarget error: Pointer 0x{{0*}}[[#HOST_ADDR]] was not present on the device upon entry to the region.
// CHECK: omptarget error: Call to targetDataBegin failed, abort target.
// CHECK: omptarget error: Failed to process data before launching the kernel.
-// CHECK: omptarget fatal error 1: failure of target construct while offloading
-// is mandatory
+// CHECK: omptarget fatal error 1: failure of target construct while offloading is mandatory
+// clang-format on
#pragma omp target map(alloc : x) map(present, alloc : x) map(tofrom : x)
{
printf("%d\n", x);
diff --git a/offload/test/mapping/map_ordering_tgt_exit_data_always_always.c b/offload/test/mapping/map_ordering_tgt_exit_data_always_always.c
index 0e00f9efc8289..f4420df483ede 100644
--- a/offload/test/mapping/map_ordering_tgt_exit_data_always_always.c
+++ b/offload/test/mapping/map_ordering_tgt_exit_data_always_always.c
@@ -19,8 +19,8 @@ int main() {
x = 222;
}
#pragma omp target exit data map(always, from : x) map(always, from : x)
- // DEBUG: omptarget --> Moving 4 bytes (tgt:0x{{.*}}) -> (hst:0x{{.*}})
- // DEBUG-NOT: omptarget --> Moving 4 bytes
+ // DEBUG: omptarget --> Moving {{.*}} bytes (tgt:0x{{.*}}) -> (hst:0x{{.*}})
+ // DEBUG-NOT: omptarget --> Moving {{.*}} bytes
}
printf("%d\n", x); // CHECK: 222
diff --git a/offload/test/mapping/map_ordering_tgt_exit_data_delete_from.c b/offload/test/mapping/map_ordering_tgt_exit_data_delete_from.c
index 896bc8bb18efa..2360ac0dd9a45 100644
--- a/offload/test/mapping/map_ordering_tgt_exit_data_delete_from.c
+++ b/offload/test/mapping/map_ordering_tgt_exit_data_delete_from.c
@@ -9,6 +9,7 @@ int main() {
#pragma omp target enter data map(alloc : x) map(to : x)
#pragma omp target map(present, alloc : x)
{
+ // NOTE: It's ok for this to be 111 under "unified_shared_memory"
printf("%d\n", x); // CHECK-NOT: 111
x = 222;
}
diff --git a/offload/test/mapping/map_ordering_tgt_exit_data_delete_from_assumedsize.c b/offload/test/mapping/map_ordering_tgt_exit_data_delete_from_assumedsize.c
index 1e4cef4b66335..e570dfbf6a7fb 100644
--- a/offload/test/mapping/map_ordering_tgt_exit_data_delete_from_assumedsize.c
+++ b/offload/test/mapping/map_ordering_tgt_exit_data_delete_from_assumedsize.c
@@ -1,6 +1,8 @@
// RUN: %libomptarget-compile-generic -fopenmp-version=60
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefix=CHECK
// RUN: env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 \
-// RUN: | %fcheck-generic -check-prefix=DEBUG -check-prefix=CHECK
+// RUN: | %fcheck-generic -check-prefix=DEBUG
// REQUIRES: libomptarget-debug
// The from on target_exit_data should result in a data-transfer of 4 bytes,
@@ -20,16 +22,17 @@ int main() {
#pragma omp target data map(alloc : x)
{
#pragma omp target enter data map(alloc : x) map(to : x)
+// DEBUG-NOT: omptarget --> Moving {{.*}} bytes (hst:0x{{.*}}) -> (tgt:0x{{.*}})
#pragma omp target map(present, alloc : x)
{
+ // NOTE: It's ok for this to be 111 under "unified_shared_memory"
printf("In tgt: %d\n", x[1]); // CHECK-NOT: In tgt: 111
x[1] = 222;
}
-// DEBUG: omptarget --> Found skipped FROM entry
-// DEBUG-SAME: HstPtr=0x[[#%x,HOST_ADDR:]] size=[[#%u,SIZE:]]
-// DEBUG-SAME: within region being deleted
-// DEBUG: omptarget --> Moving [[#SIZE]] bytes
-// DEBUG-SAME: (tgt:0x{{.*}}) -> (hst:0x{{0*}}[[#HOST_ADDR]])
+// clang-format off
+// DEBUG: omptarget --> Found skipped FROM entry: HstPtr=0x[[#%x,HOST_ADDR:]] size=[[#%u,SIZE:]] within region being deleted
+// DEBUG: omptarget --> Moving [[#SIZE]] bytes (tgt:0x{{.*}}) -> (hst:0x{{0*}}[[#HOST_ADDR]])
+// clang-format on
#pragma omp target exit data map(delete : p1x[ : ]) map(from : p2x[1])
printf("%d\n", x[1]); // CHECK: 222
}
diff --git a/offload/test/mapping/map_ordering_tgt_exit_data_from_delete_assumedsize.c b/offload/test/mapping/map_ordering_tgt_exit_data_from_delete_assumedsize.c
index d939bb66235c1..eeabed191b223 100644
--- a/offload/test/mapping/map_ordering_tgt_exit_data_from_delete_assumedsize.c
+++ b/offload/test/mapping/map_ordering_tgt_exit_data_from_delete_assumedsize.c
@@ -1,10 +1,10 @@
// RUN: %libomptarget-compile-generic -fopenmp-version=60
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefix=CHECK
// RUN: env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 \
-// RUN: | %fcheck-generic -check-prefix=DEBUG -check-prefix=CHECK
+// RUN: | %fcheck-generic -check-prefix=DEBUG
// REQUIRES: libomptarget-debug
-#include <stdio.h>
-
// The from on target_exit_data should result in a data-transfer of 4 bytes,
// even if when "delete" is honored first, and by the time "from" is
// encountered, the ref-count had already been 0 (i.e. it's not transitioning
@@ -22,16 +22,18 @@ int main() {
#pragma omp target data map(alloc : x)
{
#pragma omp target enter data map(alloc : x) map(to : x)
+// DEBUG-NOT: omptarget --> Moving {{.*}} bytes (hst:0x{{.*}}) -> (tgt:0x{{.*}})
#pragma omp target map(present, alloc : x)
{
- fprintf(stderr, "In tgt: %d\n", x[1]); // CHECK-NOT: In tgt: 111
+ // NOTE: It's ok for this to be 111 under "unified_shared_memory"
+ printf("In tgt: %d\n", x[1]); // CHECK-NOT: In tgt: 111
x[1] = 222;
}
- // DEBUG: omptarget --> Pointer HstPtr=0x[[#%x,HOST_ADDR:]] falls within a
- // DEBUG-SAME: range previously marked for deletion
- // DEBUG: omptarget --> Moving {{.*}} bytes
- // DEBUG-SAME: (tgt:0x{{.*}}) -> (hst:0x{{0*}}[[#HOST_ADDR]])
+// clang-format off
+// DEBUG: omptarget --> Pointer HstPtr=0x[[#%x,HOST_ADDR:]] falls within a range previously marked for deletion
+// DEBUG: omptarget --> Moving {{.*}} bytes (tgt:0x{{.*}}) -> (hst:0x{{0*}}[[#HOST_ADDR]])
+// clang-format on
#pragma omp target exit data map(from : p2x[0]) map(delete : p1x[ : ])
- fprintf(stderr, "%d\n", x[1]); // CHECK: 222
+ printf("%d\n", x[1]); // CHECK: 222
}
}
>From f03d646b2f4eb7bca05df47161e55854b9cac3e1 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Thu, 5 Feb 2026 14:56:09 -0800
Subject: [PATCH 14/16] Remove two tests that were useful, but not strictly
OpenMP compliant.
---
...p_ordering_tgt_alloc_from_to_assumedsize.c | 30 ------------
..._alloc_from_to_structmembers_assumedsize.c | 49 -------------------
2 files changed, 79 deletions(-)
delete mode 100644 offload/test/mapping/map_ordering_tgt_alloc_from_to_assumedsize.c
delete mode 100644 offload/test/mapping/map_ordering_tgt_alloc_from_to_structmembers_assumedsize.c
diff --git a/offload/test/mapping/map_ordering_tgt_alloc_from_to_assumedsize.c b/offload/test/mapping/map_ordering_tgt_alloc_from_to_assumedsize.c
deleted file mode 100644
index 5612aa5cb01c1..0000000000000
--- a/offload/test/mapping/map_ordering_tgt_alloc_from_to_assumedsize.c
+++ /dev/null
@@ -1,30 +0,0 @@
-// RUN: %libomptarget-compile-generic -fopenmp-version=60
-// RUN: %libomptarget-run-generic | %fcheck-generic
-
-// The test ensures that even though the from/to maps are on
-// a list-item that looks different from the "alloc" list-items,
-// which are encountered before the to/from entries at run time,
-// the transfers still happen.
-//
-// NOTE: This is not fully OpenMP compliant, but is a simpler
-// proxy for when this happens via mappers.
-
-#include <omp.h>
-#include <stdio.h>
-
-int main() {
- int x[10];
- x[1] = 111;
-
- int *xp1, *xp2;
- xp1 = xp2 = &x[0];
-
-#pragma omp target map(alloc : x[ : ]) map(from : xp1[1]) map(to : xp1[1]) \
- map(alloc : xp2[ : ])
- {
- printf("%d\n", x[1]); // CHECK: 111
- x[1] = x[1] + 111;
- }
-
- printf("%d\n", x[1]); // CHECK: 222
-}
diff --git a/offload/test/mapping/map_ordering_tgt_alloc_from_to_structmembers_assumedsize.c b/offload/test/mapping/map_ordering_tgt_alloc_from_to_structmembers_assumedsize.c
deleted file mode 100644
index 54fa075dc4827..0000000000000
--- a/offload/test/mapping/map_ordering_tgt_alloc_from_to_structmembers_assumedsize.c
+++ /dev/null
@@ -1,49 +0,0 @@
-// RUN: %libomptarget-compile-generic -fopenmp-version=60
-// RUN: %libomptarget-run-generic 2>&1 \
-// RUN: | %fcheck-generic -check-prefix=CHECK
-// RUN: env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 \
-// RUN: | %fcheck-generic -check-prefix=DEBUG
-// REQUIRES: libomptarget-debug
-
-#include <omp.h>
-#include <stdio.h>
-
-// This test shows that when a "release" is being handled, we look at
-// all previously skipped data-transfers within the range of the memory being
-// released. Here, when "s[:]" is being released, we honor the two "from"s
-// on sp1->x and sp1->y.
-
-typedef struct {
- int x;
- int y;
-} S;
-
-int main() {
- S s[10];
- s[1].x = 111;
- s[1].y = 111;
-
- S *sp1 = &s[1];
-
-// clang-format off
-// DEBUG: omptarget --> HstPtrBegin 0x[[#%x,HOST_ADDRX:]] was newly allocated for the current region
-// DEBUG: omptarget --> Moving [[#%u,SIZEX:]] bytes (hst:0x{{0*}}[[#HOST_ADDRX]]) -> (tgt:0x{{.*}})
-// DEBUG: omptarget --> HstPtrBegin 0x[[#%x,HOST_ADDRY:]] was newly allocated for the current region
-// DEBUG: omptarget --> Moving [[#%u,SIZEY:]] bytes (hst:0x{{0*}}[[#HOST_ADDRY]]) -> (tgt:0x{{.*}})
-// clang-format on
-#pragma omp target map(alloc : s[ : ]) map(from : sp1 -> x, sp1->y) \
- map(to : sp1->x, sp1->y) map(alloc : s[ : ])
- {
- fprintf(stderr, "%d %d\n", s[1].x, s[1].y); // CHECK: 111 111
- s[1].x = s[1].x + 111;
- s[1].y = s[1].y + 111;
- }
-
- // clang-format off
- // DEBUG: omptarget --> Found skipped FROM entry: HstPtr=0x{{0*}}[[#HOST_ADDRY]] size=[[#SIZEY]] within region being deleted
- // DEBUG: omptarget --> Moving [[#SIZEY]] bytes (tgt:0x{{.*}}) -> (hst:0x{{0*}}[[#HOST_ADDRY]])
- // DEBUG: omptarget --> Found skipped FROM entry: HstPtr=0x{{0*}}[[#HOST_ADDRX]] size=[[#SIZEX]] within region being deleted
- // DEBUG: omptarget --> Moving [[#SIZEX]] bytes (tgt:0x{{.*}}) -> (hst:0x{{0*}}[[#HOST_ADDRX]])
- // clang-format on
- fprintf(stderr, "%d %d\n", s[1].x, s[1].y); // CHECK: 222 222
-}
>From 115e36d4c83909a8ace1e5ce4f0717a51683f855 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Thu, 5 Feb 2026 15:15:26 -0800
Subject: [PATCH 15/16] Clang-format fixes.
---
offload/include/OpenMP/Mapping.h | 15 ++++++------
offload/libomptarget/omptarget.cpp | 23 +++++++++++--------
...ring_ptee_tgt_alloc_mapper_alloc_from_to.c | 19 +++++++--------
..._alloc_tgt_mapper_present_delete_from_to.c | 13 +++++------
.../map_ordering_tgt_alloc_present_tofrom.c | 15 ++++++------
...ng_tgt_exit_data_delete_from_assumedsize.c | 10 ++++----
...ng_tgt_exit_data_from_delete_assumedsize.c | 10 ++++----
7 files changed, 57 insertions(+), 48 deletions(-)
diff --git a/offload/include/OpenMP/Mapping.h b/offload/include/OpenMP/Mapping.h
index b052504b425e3..88f6d45accb5f 100644
--- a/offload/include/OpenMP/Mapping.h
+++ b/offload/include/OpenMP/Mapping.h
@@ -528,14 +528,13 @@ struct StateInfoTy {
/// Check if a pointer falls within any of the newly allocated ranges.
/// Returns true if the pointer is within a newly allocated region.
bool wasNewlyAllocated(void *Ptr) const {
- return llvm::any_of(
- NewAllocations, [&](const auto &Alloc) {
- void *AllocPtr = Alloc.first;
- int64_t AllocSize = Alloc.second;
- return Ptr >= AllocPtr &&
- Ptr < static_cast<void *>(static_cast<char *>(AllocPtr) +
- AllocSize);
- });
+ return llvm::any_of(NewAllocations, [&](const auto &Alloc) {
+ void *AllocPtr = Alloc.first;
+ int64_t AllocSize = Alloc.second;
+ return Ptr >= AllocPtr &&
+ Ptr <
+ static_cast<void *>(static_cast<char *>(AllocPtr) + AllocSize);
+ });
}
};
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index b5b4cf2d984b0..366a268eef98d 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -1246,11 +1246,13 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
void *DeletePtr = DeleteEntry.first;
int64_t DeleteSize = DeleteEntry.second;
if (HstPtrBegin >= DeletePtr &&
- HstPtrBegin < static_cast<void *>(static_cast<char *>(DeletePtr) + DeleteSize)) {
+ HstPtrBegin < static_cast<void *>(static_cast<char *>(DeletePtr) +
+ DeleteSize)) {
ODBG(ODT_Mapping)
<< "Pointer HstPtr=" << HstPtrBegin
<< " falls within a range previously marked for deletion ["
- << DeletePtr << ", " << static_cast<char *>(DeletePtr) + DeleteSize
+ << DeletePtr << ", "
+ << static_cast<char *>(DeletePtr) + DeleteSize
<< ") with size=" << DeleteSize;
return true;
}
@@ -1280,8 +1282,9 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
TPR.getEntry());
if (Ret != OFFLOAD_SUCCESS)
return OFFLOAD_FAIL;
- } else if (!FromCopyBackAlreadyDone && (IsMapFromOnNonHostNonZeroData &&
- !IsLastOrHasAlwaysOrWasForceDeleted())) {
+ } else if (!FromCopyBackAlreadyDone &&
+ (IsMapFromOnNonHostNonZeroData &&
+ !IsLastOrHasAlwaysOrWasForceDeleted())) {
// We can have cases like the following:
// p1 = p2 = &x;
// ... map(storage: p1[:]) map(from: p2[1:1])
@@ -1327,16 +1330,18 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
ODBG(ODT_Mapping)
<< "Found skipped FROM entry: HstPtr=" << FromBeginPtr
<< " size=" << FromDataSize << " within region being deleted ["
- << reinterpret_cast<void*>(DeleteBeginPtrInt) << ", " << reinterpret_cast<void*>(DeleteEndPtrInt) << ")";
+ << reinterpret_cast<void *>(DeleteBeginPtrInt) << ", "
+ << reinterpret_cast<void *>(DeleteEndPtrInt) << ")";
// Calculate offset within the target pointer
int64_t Offset = FromBeginPtrInt - DeleteBeginPtrInt;
- void *FromTgtBeginPtr = static_cast<void *>(static_cast<char *>(TgtPtrBegin) + Offset);
+ void *FromTgtBeginPtr =
+ static_cast<void *>(static_cast<char *>(TgtPtrBegin) + Offset);
// Perform the retrieval for this skipped entry
- int Ret =
- PerformFromRetrieval(reinterpret_cast<void *>(FromBeginPtrInt), FromTgtBeginPtr,
- FromDataSize, TPR.getEntry());
+ int Ret = PerformFromRetrieval(
+ reinterpret_cast<void *>(FromBeginPtrInt), FromTgtBeginPtr,
+ FromDataSize, TPR.getEntry());
if (Ret != OFFLOAD_SUCCESS)
return OFFLOAD_FAIL;
diff --git a/offload/test/mapping/map_ordering_ptee_tgt_alloc_mapper_alloc_from_to.c b/offload/test/mapping/map_ordering_ptee_tgt_alloc_mapper_alloc_from_to.c
index 9a94cab804b71..636fb2d87507f 100644
--- a/offload/test/mapping/map_ordering_ptee_tgt_alloc_mapper_alloc_from_to.c
+++ b/offload/test/mapping/map_ordering_ptee_tgt_alloc_mapper_alloc_from_to.c
@@ -19,7 +19,8 @@ typedef struct {
int *q;
} S;
#pragma omp declare mapper(my_mapper : S s) map(alloc : s.p, s.p[0 : 10]) \
- map(from : s.p[0 : 10]) map(to : s.p[0 : 10]) map(alloc : s.p[0 : 10]) // (2)
+ map(from : s.p[0 : 10]) map(to : s.p[0 : 10]) \
+ map(alloc : s.p[0 : 10]) // (2)
S s1;
int main() {
@@ -27,10 +28,10 @@ int main() {
x[1] = 111;
s1.q = s1.p = &x[0];
-// clang-format off
-// DEBUG: omptarget --> HstPtrBegin 0x[[#%x,HOST_ADDRX:]] was newly allocated for the current region
-// DEBUG: omptarget --> Moving [[#%u,SIZEX:]] bytes (hst:0x{{0*}}[[#HOST_ADDRX]]) -> (tgt:0x{{.*}})
-// clang-format on
+ // clang-format off
+ // DEBUG: omptarget --> HstPtrBegin 0x[[#%x,HOST_ADDRX:]] was newly allocated for the current region
+ // DEBUG: omptarget --> Moving [[#%u,SIZEX:]] bytes (hst:0x{{0*}}[[#HOST_ADDRX]]) -> (tgt:0x{{.*}})
+ // clang-format on
#pragma omp target map(alloc : s1.p[0 : 10]) \
map(mapper(my_mapper), tofrom : s1) // (1)
{
@@ -38,9 +39,9 @@ int main() {
s1.p[1] = s1.p[1] + 111;
}
-// clang-format off
-// DEBUG: omptarget --> Found skipped FROM entry: HstPtr=0x{{0*}}[[#HOST_ADDRX]] size=[[#SIZEX]] within region being deleted
-// DEBUG: omptarget --> Moving [[#SIZEX]] bytes (tgt:0x{{.*}}) -> (hst:0x{{0*}}[[#HOST_ADDRX]])
-// clang-format on
+ // clang-format off
+ // DEBUG: omptarget --> Found skipped FROM entry: HstPtr=0x{{0*}}[[#HOST_ADDRX]] size=[[#SIZEX]] within region being deleted
+ // DEBUG: omptarget --> Moving [[#SIZEX]] bytes (tgt:0x{{.*}}) -> (hst:0x{{0*}}[[#HOST_ADDRX]])
+ // clang-format on
printf("%d\n", s1.p[1]); // CHECK: 222
}
diff --git a/offload/test/mapping/map_ordering_ptee_tgt_data_alloc_tgt_mapper_present_delete_from_to.c b/offload/test/mapping/map_ordering_ptee_tgt_data_alloc_tgt_mapper_present_delete_from_to.c
index 791ff3ddb2a3b..ceef9a5848632 100644
--- a/offload/test/mapping/map_ordering_ptee_tgt_data_alloc_tgt_mapper_present_delete_from_to.c
+++ b/offload/test/mapping/map_ordering_ptee_tgt_data_alloc_tgt_mapper_present_delete_from_to.c
@@ -11,9 +11,8 @@
// However, there should be a "from" transfer at the end of (2) because of the
// "delete" on the mapper.
-// FIXME: This currently fails, but should start passing once ATTACH-style maps are
-// enabled for mappers (#166874).
-// UNSUPPORTED: true
+// FIXME: This currently fails, but should start passing once ATTACH-style maps
+// are enabled for mappers (#166874). UNSUPPORTED: true
#include <stdio.h>
@@ -42,8 +41,8 @@ int main() {
}
printf("%d\n", s1.p[1]); // CHECK: 222
}
-// clang-format off
-// DEBUG: omptarget --> Found skipped FROM entry: HstPtr=0x[[#%x,HOST_ADDR:]] size=[[#%u,SIZE:]] within region being deleted
-// DEBUG: omptarget --> Moving [[#SIZE]] bytes (tgt:0x{{.*}}) -> (hst:0x{{0*}}[[#HOST_ADDR]])
-// clang-format on
+ // clang-format off
+ // DEBUG: omptarget --> Found skipped FROM entry: HstPtr=0x[[#%x,HOST_ADDR:]] size=[[#%u,SIZE:]] within region being deleted
+ // DEBUG: omptarget --> Moving [[#SIZE]] bytes (tgt:0x{{.*}}) -> (hst:0x{{0*}}[[#HOST_ADDR]])
+ // clang-format on
}
diff --git a/offload/test/mapping/map_ordering_tgt_alloc_present_tofrom.c b/offload/test/mapping/map_ordering_tgt_alloc_present_tofrom.c
index 32d5cf4b1b06b..e3725bb8967da 100644
--- a/offload/test/mapping/map_ordering_tgt_alloc_present_tofrom.c
+++ b/offload/test/mapping/map_ordering_tgt_alloc_present_tofrom.c
@@ -8,13 +8,14 @@ int main() {
// CHECK: addr=0x[[#%x,HOST_ADDR:]], size=[[#%u,SIZE:]]
int x = 111;
fprintf(stderr, "addr=%p, size=%ld\n", &x, sizeof(x));
-// clang-format off
-// CHECK: omptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes)
-// CHECK: omptarget error: Pointer 0x{{0*}}[[#HOST_ADDR]] was not present on the device upon entry to the region.
-// CHECK: omptarget error: Call to targetDataBegin failed, abort target.
-// CHECK: omptarget error: Failed to process data before launching the kernel.
-// CHECK: omptarget fatal error 1: failure of target construct while offloading is mandatory
-// clang-format on
+
+ // clang-format off
+ // CHECK: omptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes)
+ // CHECK: omptarget error: Pointer 0x{{0*}}[[#HOST_ADDR]] was not present on the device upon entry to the region.
+ // CHECK: omptarget error: Call to targetDataBegin failed, abort target.
+ // CHECK: omptarget error: Failed to process data before launching the kernel.
+ // CHECK: omptarget fatal error 1: failure of target construct while offloading is mandatory
+ // clang-format on
#pragma omp target map(alloc : x) map(present, alloc : x) map(tofrom : x)
{
printf("%d\n", x);
diff --git a/offload/test/mapping/map_ordering_tgt_exit_data_delete_from_assumedsize.c b/offload/test/mapping/map_ordering_tgt_exit_data_delete_from_assumedsize.c
index e570dfbf6a7fb..1c266ff4bbfa0 100644
--- a/offload/test/mapping/map_ordering_tgt_exit_data_delete_from_assumedsize.c
+++ b/offload/test/mapping/map_ordering_tgt_exit_data_delete_from_assumedsize.c
@@ -29,11 +29,13 @@ int main() {
printf("In tgt: %d\n", x[1]); // CHECK-NOT: In tgt: 111
x[1] = 222;
}
-// clang-format off
-// DEBUG: omptarget --> Found skipped FROM entry: HstPtr=0x[[#%x,HOST_ADDR:]] size=[[#%u,SIZE:]] within region being deleted
-// DEBUG: omptarget --> Moving [[#SIZE]] bytes (tgt:0x{{.*}}) -> (hst:0x{{0*}}[[#HOST_ADDR]])
-// clang-format on
+
#pragma omp target exit data map(delete : p1x[ : ]) map(from : p2x[1])
+ // clang-format off
+ // DEBUG: omptarget --> Found skipped FROM entry: HstPtr=0x[[#%x,HOST_ADDR:]] size=[[#%u,SIZE:]] within region being deleted
+ // DEBUG: omptarget --> Moving [[#SIZE]] bytes (tgt:0x{{.*}}) -> (hst:0x{{0*}}[[#HOST_ADDR]])
+ // clang-format on
+
printf("%d\n", x[1]); // CHECK: 222
}
}
diff --git a/offload/test/mapping/map_ordering_tgt_exit_data_from_delete_assumedsize.c b/offload/test/mapping/map_ordering_tgt_exit_data_from_delete_assumedsize.c
index eeabed191b223..963e709411c92 100644
--- a/offload/test/mapping/map_ordering_tgt_exit_data_from_delete_assumedsize.c
+++ b/offload/test/mapping/map_ordering_tgt_exit_data_from_delete_assumedsize.c
@@ -29,11 +29,13 @@ int main() {
printf("In tgt: %d\n", x[1]); // CHECK-NOT: In tgt: 111
x[1] = 222;
}
-// clang-format off
-// DEBUG: omptarget --> Pointer HstPtr=0x[[#%x,HOST_ADDR:]] falls within a range previously marked for deletion
-// DEBUG: omptarget --> Moving {{.*}} bytes (tgt:0x{{.*}}) -> (hst:0x{{0*}}[[#HOST_ADDR]])
-// clang-format on
+
#pragma omp target exit data map(from : p2x[0]) map(delete : p1x[ : ])
+ // clang-format off
+ // DEBUG: omptarget --> Pointer HstPtr=0x[[#%x,HOST_ADDR:]] falls within a range previously marked for deletion
+ // DEBUG: omptarget --> Moving {{.*}} bytes (tgt:0x{{.*}}) -> (hst:0x{{0*}}[[#HOST_ADDR]])
+ // clang-format on
+
printf("%d\n", x[1]); // CHECK: 222
}
}
>From a68d2be88098fa7ba6f8e0fed91b1a83b6d153e4 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Tue, 10 Feb 2026 06:38:30 -0800
Subject: [PATCH 16/16] Make from+from overlap case more robust.
Also improve robustness of multiple FROM entries of different sizes
but same starting address.
---
offload/include/OpenMP/Mapping.h | 91 +++++++++--
offload/libomptarget/OpenMP/Mapping.cpp | 3 +-
offload/libomptarget/omptarget.cpp | 148 ++++++++++--------
...ring_ptee_tgt_alloc_mapper_alloc_from_to.c | 2 +-
...ng_tgt_exit_data_delete_from_assumedsize.c | 2 +-
...ng_tgt_exit_data_from_delete_assumedsize.c | 2 +-
...dering_tgt_exit_data_from_mapper_overlap.c | 49 ++++++
7 files changed, 211 insertions(+), 86 deletions(-)
create mode 100644 offload/test/mapping/map_ordering_tgt_exit_data_from_mapper_overlap.c
diff --git a/offload/include/OpenMP/Mapping.h b/offload/include/OpenMP/Mapping.h
index 88f6d45accb5f..e4024abf26690 100644
--- a/offload/include/OpenMP/Mapping.h
+++ b/offload/include/OpenMP/Mapping.h
@@ -512,12 +512,16 @@ struct StateInfoTy {
llvm::DenseMap<void *, int64_t> SkippedFromEntries;
/// Host pointers for which we have triggered a FROM transfer at some point
- /// during targetDataEnd. Used to avoid duplicate transfers.
- llvm::SmallSet<void *, 32> TransferredFromPtrs;
+ /// during targetDataEnd. It's used to avoid duplicate transfers.
+ /// Key: host pointer, Value: transferred size.
+ llvm::DenseMap<void *, int64_t> TransferredFromEntries;
- /// Starting host address and size of the DELETE entries previouly processed
- /// for the current region. Key: host pointer, Value: allocated size.
- llvm::DenseMap<void *, int64_t> DeleteEntries;
+ /// Starting host address and size of entries whose ref-count went to zero.
+ /// This includes entries released through explicit DELETE, or normal
+ /// ref-count decrements. It's used to ensure transfers are performed for FROM
+ /// entries whose ref-count is already zero when the entry is encountered.
+ /// Key: host pointer, Value: size.
+ llvm::DenseMap<void *, int64_t> ReleasedEntries;
StateInfoTy() = default;
@@ -525,16 +529,75 @@ struct StateInfoTy {
StateInfoTy(const StateInfoTy &) = delete;
StateInfoTy &operator=(const StateInfoTy &) = delete;
+private:
+ /// Helper to find an entry in \p EntryMap that contains the pointer.
+ /// Returns the matching entry if found, otherwise std::nullopt.
+ std::optional<std::pair<void *, int64_t>>
+ findEntryForPtr(void *Ptr,
+ const llvm::DenseMap<void *, int64_t> &EntryMap) const {
+ for (const auto &Entry : EntryMap) {
+ void *EntryBegin = Entry.first;
+ int64_t EntrySize = Entry.second;
+ if (Ptr >= EntryBegin &&
+ Ptr < static_cast<void *>(static_cast<char *>(EntryBegin) +
+ EntrySize)) {
+ return Entry;
+ }
+ }
+ return std::nullopt;
+ }
+
+public:
/// Check if a pointer falls within any of the newly allocated ranges.
- /// Returns true if the pointer is within a newly allocated region.
- bool wasNewlyAllocated(void *Ptr) const {
- return llvm::any_of(NewAllocations, [&](const auto &Alloc) {
- void *AllocPtr = Alloc.first;
- int64_t AllocSize = Alloc.second;
- return Ptr >= AllocPtr &&
- Ptr <
- static_cast<void *>(static_cast<char *>(AllocPtr) + AllocSize);
- });
+ /// Returns the matching entry if found, otherwise std::nullopt.
+ std::optional<std::pair<void *, int64_t>> wasNewlyAllocated(void *Ptr) const {
+ return findEntryForPtr(Ptr, NewAllocations);
+ }
+
+ /// Check if a pointer range [Ptr, Ptr+Size) is fully contained within any
+ /// previously completed FROM transfer.
+ /// Returns the matching entry if found, otherwise std::nullopt.
+ std::optional<std::pair<void *, int64_t>>
+ wasTransferredFrom(void *Ptr, int64_t Size) const {
+ uintptr_t CheckBegin = reinterpret_cast<uintptr_t>(Ptr);
+ uintptr_t CheckEnd = CheckBegin + Size;
+
+ for (const auto &Entry : TransferredFromEntries) {
+ void *RangePtr = Entry.first;
+ int64_t RangeSize = Entry.second;
+ uintptr_t RangeBegin = reinterpret_cast<uintptr_t>(RangePtr);
+ uintptr_t RangeEnd = RangeBegin + RangeSize;
+
+ if (CheckBegin >= RangeBegin && CheckEnd <= RangeEnd) {
+ return Entry;
+ }
+ }
+ return std::nullopt;
+ }
+
+ /// Check if a pointer falls within any released entry's range.
+ /// Returns the matching entry if found, otherwise std::nullopt.
+ std::optional<std::pair<void *, int64_t>>
+ wasPreviouslyReleased(void *Ptr) const {
+ return findEntryForPtr(Ptr, ReleasedEntries);
+ }
+
+ /// Add a skipped FROM entry. Only updates the entry if this is a new pointer
+ /// or if the new size is larger than the existing entry.
+ void addSkippedFromEntry(void *Ptr, int64_t Size) {
+ auto It = SkippedFromEntries.find(Ptr);
+ if (It == SkippedFromEntries.end() || Size > It->second) {
+ SkippedFromEntries[Ptr] = Size;
+ }
+ }
+
+ /// Add a transferred FROM entry. Only updates the entry if this is a new
+ /// pointer or if the new size is larger than the existing entry.
+ void addTransferredFromEntry(void *Ptr, int64_t Size) {
+ auto It = TransferredFromEntries.find(Ptr);
+ if (It == TransferredFromEntries.end() || Size > It->second) {
+ TransferredFromEntries[Ptr] = Size;
+ }
}
};
diff --git a/offload/libomptarget/OpenMP/Mapping.cpp b/offload/libomptarget/OpenMP/Mapping.cpp
index 6f075f6545b3d..1bb2e424bd083 100644
--- a/offload/libomptarget/OpenMP/Mapping.cpp
+++ b/offload/libomptarget/OpenMP/Mapping.cpp
@@ -345,7 +345,8 @@ TargetPointerResultTy MappingInfoTy::getTargetPointer(
auto WasNewlyAllocatedForCurrentRegion = [&]() {
if (!StateInfo)
return false;
- bool WasNewlyAllocated = StateInfo->wasNewlyAllocated(HstPtrBegin);
+ bool WasNewlyAllocated =
+ StateInfo->wasNewlyAllocated(HstPtrBegin).has_value();
if (WasNewlyAllocated)
ODBG(ODT_Mapping) << "HstPtrBegin " << HstPtrBegin
<< " was newly allocated for the current region";
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index 366a268eef98d..344c388e794af 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -672,7 +672,7 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
<< ").";
return OFFLOAD_FAIL;
} else if (TgtPtrBegin && HasPresentModifier &&
- StateInfo->wasNewlyAllocated(HstPtrBegin)) {
+ StateInfo->wasNewlyAllocated(HstPtrBegin).has_value()) {
// For "PRESENT" entries, we may have cases like the following:
// int *xp = &x[0];
// map(alloc: x[:]) map(present, alloc: xp[1])
@@ -860,11 +860,11 @@ int processAttachEntries(DeviceTy &Device, StateInfoTy &StateInfo,
// Lambda to check if a pointer was newly allocated
auto WasNewlyAllocated = [&](void *Ptr, const char *PtrName) {
- bool IsNewlyAllocated = StateInfo.wasNewlyAllocated(Ptr);
+ bool WasNewlyAllocated = StateInfo.wasNewlyAllocated(Ptr).has_value();
ODBG(ODT_Mapping) << "Attach " << PtrName << " " << Ptr
<< " was newly allocated: "
- << (IsNewlyAllocated ? "yes" : "no");
- return IsNewlyAllocated;
+ << (WasNewlyAllocated ? "yes" : "no");
+ return WasNewlyAllocated;
};
// Only process ATTACH if either the pointee or the pointer was newly
@@ -1173,18 +1173,21 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
if (!TPR.isPresent())
continue;
- // Track force-deleted pointers so we can use this information if we
- // encounter FROM entries for the same pointer later on.
- if (ForceDelete && TPR.Flags.IsLast) {
+ // Track entries whose ref-count went to zero (IsLast=true) so that we
+ // can honor any subsequently encountered FROM entries that fall within
+ // their range.
+ if (TPR.Flags.IsLast) {
// For assumed-size arrays like map(delete: p[:]), the compiler provides
// no size information, so we need to get the actual allocated extent from
// the HDTT entry.
- int64_t AllocatedSize =
+ void *ReleasedHstPtrBegin =
+ reinterpret_cast<void *>(TPR.getEntry()->HstPtrBegin);
+ int64_t ReleasedSize =
TPR.getEntry()->HstPtrEnd - TPR.getEntry()->HstPtrBegin;
- ODBG(ODT_Mapping) << "Marking HstPtr=" << HstPtrBegin
- << " for deletion with allocated size="
- << AllocatedSize;
- StateInfo->DeleteEntries[HstPtrBegin] = AllocatedSize;
+ ODBG(ODT_Mapping) << "Tracking released entry: HstPtr="
+ << ReleasedHstPtrBegin << ", Size=" << ReleasedSize
+ << ", ForceDelete=" << ForceDelete;
+ StateInfo->ReleasedEntries[ReleasedHstPtrBegin] = ReleasedSize;
}
// Move data back to the host
@@ -1194,6 +1197,28 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
// Lambda to perform the actual FROM data retrieval from device to host
auto PerformFromRetrieval = [&](void *HstPtr, void *TgtPtr, int64_t Size,
HostDataToTargetTy *Entry) -> int {
+ // Check if this FROM transfer can be skipped.
+ //
+ // This is an optimization that may help in rare cases when we have
+ // multiple overlapping FROM entries. e.g.
+ //
+ // ... map(always, from: x) map(always, from: x)
+ // ... map(delete: x) map(from: x) map(from: x)
+ //
+ // If we think the overhead makes it not worh it, we can remove it.
+ if (auto TransferredEntry = StateInfo->wasTransferredFrom(HstPtr, Size)) {
+ void *TransferredPtr = TransferredEntry->first;
+ int64_t TransferredSize = TransferredEntry->second;
+ ODBG(ODT_Mapping) << "FROM entry HstPtr=" << HstPtr << " size=" << Size
+ << " already transferred within [" << TransferredPtr
+ << ", "
+ << static_cast<void *>(
+ static_cast<char *>(TransferredPtr) +
+ TransferredSize)
+ << ")";
+ return OFFLOAD_SUCCESS;
+ }
+
ODBG(ODT_Mapping) << "Moving " << Size << " bytes (tgt:" << TgtPtr
<< ") -> (hst:" << HstPtr << ")";
TIMESCOPE_WITH_DETAILS_AND_IDENT(
@@ -1222,13 +1247,14 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
return OFFLOAD_FAIL;
}
+ // Track this transfer to avoid duplicate transfers later on.
+ StateInfo->addTransferredFromEntry(HstPtr, Size);
+
return OFFLOAD_SUCCESS;
};
- // Lambda to check if this pointer was previously marked for deletion.
- // Such a pointer would have had "IsLast" set to true when its DELETE entry
- // was processed. So, the flag wouldn't be set for any FROM entries seen
- // later on.
+ // Lambda to check if this pointer was previously released.
+ //
// This is needed to handle cases like the following:
// p1 = p2 = &x;
// ... map(delete: p1[:]) map(from: p2[0:1])
@@ -1240,51 +1266,35 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
// the always-modifier or delete-modifier is specified, and if the map
// type is from, the original list item is updated as if the list item
// appeared in a from clause on a target_update directive.
- auto WasPreviouslyMarkedForDeletion = [&]() -> bool {
- // Check if this pointer falls within the range of any deleted entry
- for (const auto &DeleteEntry : StateInfo->DeleteEntries) {
- void *DeletePtr = DeleteEntry.first;
- int64_t DeleteSize = DeleteEntry.second;
- if (HstPtrBegin >= DeletePtr &&
- HstPtrBegin < static_cast<void *>(static_cast<char *>(DeletePtr) +
- DeleteSize)) {
- ODBG(ODT_Mapping)
- << "Pointer HstPtr=" << HstPtrBegin
- << " falls within a range previously marked for deletion ["
- << DeletePtr << ", "
- << static_cast<char *>(DeletePtr) + DeleteSize
- << ") with size=" << DeleteSize;
- return true;
- }
- }
- return false;
+ auto WasPreviouslyReleased = [&]() -> bool {
+ auto ReleasedEntry = StateInfo->wasPreviouslyReleased(HstPtrBegin);
+ if (!ReleasedEntry)
+ return false;
+
+ void *ReleasedPtr = ReleasedEntry->first;
+ int64_t ReleasedSize = ReleasedEntry->second;
+ ODBG(ODT_Mapping) << "Pointer HstPtr=" << HstPtrBegin
+ << " falls within a range previously released ["
+ << ReleasedPtr << ", "
+ << static_cast<void *>(
+ static_cast<char *>(ReleasedPtr) + ReleasedSize)
+ << ") with size=" << ReleasedSize;
+ return true;
};
- bool FromCopyBackAlreadyDone =
- StateInfo->TransferredFromPtrs.contains(HstPtrBegin);
bool IsMapFromOnNonHostNonZeroData =
HasFrom && !TPR.Flags.IsHostPointer && DataSize != 0;
- auto IsLastOrHasAlwaysOrWasForceDeleted = [&]() {
- return TPR.Flags.IsLast || HasAlways || WasPreviouslyMarkedForDeletion();
- };
- if (!FromCopyBackAlreadyDone && (IsMapFromOnNonHostNonZeroData &&
- IsLastOrHasAlwaysOrWasForceDeleted())) {
- // Track that we're doing a FROM transfer for this pointer
- // NOTE: If we don't care about the case of multiple different maps with
- // from, always, or multiple map(from)s seen after a map(delete), e.g.
- // ... map(always, from: x) map(always, from: x)
- // ... map(delete: x) map(from: x) map(from: x)
- // Then we can forego tacking TransferredFromPtrs.
- StateInfo->TransferredFromPtrs.insert(HstPtrBegin);
+ auto IsLastOrHasAlwaysOrWasReleased = [&]() {
+ return TPR.Flags.IsLast || HasAlways || WasPreviouslyReleased();
+ };
+ if (IsMapFromOnNonHostNonZeroData && IsLastOrHasAlwaysOrWasReleased()) {
Ret = PerformFromRetrieval(HstPtrBegin, TgtPtrBegin, DataSize,
TPR.getEntry());
if (Ret != OFFLOAD_SUCCESS)
return OFFLOAD_FAIL;
- } else if (!FromCopyBackAlreadyDone &&
- (IsMapFromOnNonHostNonZeroData &&
- !IsLastOrHasAlwaysOrWasForceDeleted())) {
+ } else if (IsMapFromOnNonHostNonZeroData) {
// We can have cases like the following:
// p1 = p2 = &x;
// ... map(storage: p1[:]) map(from: p2[1:1])
@@ -1307,13 +1317,18 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
// list item, one is the containing structure of the other, at least one
// is an assumed-size array, or at least one is implicitly mapped due to
// the list item also appearing in a use_device_addr clause.
- StateInfo->SkippedFromEntries[HstPtrBegin] = DataSize;
+ StateInfo->addSkippedFromEntry(HstPtrBegin, DataSize);
ODBG(ODT_Mapping) << "Skipping FROM map transfer for HstPtr="
- << HstPtrBegin << " size=" << DataSize;
- } else if (!FromCopyBackAlreadyDone && TPR.Flags.IsLast) {
- // Even if this is not a FROM entry, if the ref-count went to
- // zero (IsLast=true), we should perform any previously skipped FROM
- // transfers that fall within this entry's range.
+ << HstPtrBegin << " size=" << DataSize
+ << " (IsLast=" << TPR.Flags.IsLast << ", TotalRefCount="
+ << TPR.getEntry()->getTotalRefCount() << ")";
+ }
+
+ // If the ref-count went to zero (IsLast=true), check if any previously
+ // skipped FROM entries fall within this released entry's range.
+ if (TPR.Flags.IsLast && !StateInfo->SkippedFromEntries.empty()) {
+ uintptr_t ReleasedBeginPtrInt = TPR.getEntry()->HstPtrBegin;
+ uintptr_t ReleasedEndPtrInt = TPR.getEntry()->HstPtrEnd;
SmallVector<void *, 32> ToRemove;
for (auto &SkippedFromEntry : StateInfo->SkippedFromEntries) {
@@ -1321,20 +1336,18 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
int64_t FromDataSize = SkippedFromEntry.second;
uintptr_t FromBeginPtrInt = reinterpret_cast<uintptr_t>(FromBeginPtr);
- uintptr_t DeleteBeginPtrInt = TPR.getEntry()->HstPtrBegin;
- uintptr_t DeleteEndPtrInt = TPR.getEntry()->HstPtrEnd;
-
- // Check if skipped entry overlaps or is contained within current entry
- if (FromBeginPtrInt >= DeleteBeginPtrInt &&
- FromBeginPtrInt < DeleteEndPtrInt) {
+ // Check if this skipped FROM entry's starting pointer falls within this
+ // released entry
+ if (FromBeginPtrInt >= ReleasedBeginPtrInt &&
+ FromBeginPtrInt < ReleasedEndPtrInt) {
ODBG(ODT_Mapping)
<< "Found skipped FROM entry: HstPtr=" << FromBeginPtr
- << " size=" << FromDataSize << " within region being deleted ["
- << reinterpret_cast<void *>(DeleteBeginPtrInt) << ", "
- << reinterpret_cast<void *>(DeleteEndPtrInt) << ")";
+ << " size=" << FromDataSize << " within region being released ["
+ << reinterpret_cast<void *>(ReleasedBeginPtrInt) << ", "
+ << reinterpret_cast<void *>(ReleasedEndPtrInt) << ")";
// Calculate offset within the target pointer
- int64_t Offset = FromBeginPtrInt - DeleteBeginPtrInt;
+ int64_t Offset = FromBeginPtrInt - ReleasedBeginPtrInt;
void *FromTgtBeginPtr =
static_cast<void *>(static_cast<char *>(TgtPtrBegin) + Offset);
@@ -1345,7 +1358,6 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
if (Ret != OFFLOAD_SUCCESS)
return OFFLOAD_FAIL;
- StateInfo->TransferredFromPtrs.insert(FromBeginPtr);
ToRemove.push_back(FromBeginPtr);
}
}
diff --git a/offload/test/mapping/map_ordering_ptee_tgt_alloc_mapper_alloc_from_to.c b/offload/test/mapping/map_ordering_ptee_tgt_alloc_mapper_alloc_from_to.c
index 636fb2d87507f..5b863e05f7e43 100644
--- a/offload/test/mapping/map_ordering_ptee_tgt_alloc_mapper_alloc_from_to.c
+++ b/offload/test/mapping/map_ordering_ptee_tgt_alloc_mapper_alloc_from_to.c
@@ -40,7 +40,7 @@ int main() {
}
// clang-format off
- // DEBUG: omptarget --> Found skipped FROM entry: HstPtr=0x{{0*}}[[#HOST_ADDRX]] size=[[#SIZEX]] within region being deleted
+ // DEBUG: omptarget --> Found skipped FROM entry: HstPtr=0x{{0*}}[[#HOST_ADDRX]] size=[[#SIZEX]] within region being released
// DEBUG: omptarget --> Moving [[#SIZEX]] bytes (tgt:0x{{.*}}) -> (hst:0x{{0*}}[[#HOST_ADDRX]])
// clang-format on
printf("%d\n", s1.p[1]); // CHECK: 222
diff --git a/offload/test/mapping/map_ordering_tgt_exit_data_delete_from_assumedsize.c b/offload/test/mapping/map_ordering_tgt_exit_data_delete_from_assumedsize.c
index 1c266ff4bbfa0..5d01d29e8d47b 100644
--- a/offload/test/mapping/map_ordering_tgt_exit_data_delete_from_assumedsize.c
+++ b/offload/test/mapping/map_ordering_tgt_exit_data_delete_from_assumedsize.c
@@ -32,7 +32,7 @@ int main() {
#pragma omp target exit data map(delete : p1x[ : ]) map(from : p2x[1])
// clang-format off
- // DEBUG: omptarget --> Found skipped FROM entry: HstPtr=0x[[#%x,HOST_ADDR:]] size=[[#%u,SIZE:]] within region being deleted
+ // DEBUG: omptarget --> Found skipped FROM entry: HstPtr=0x[[#%x,HOST_ADDR:]] size=[[#%u,SIZE:]] within region being released
// DEBUG: omptarget --> Moving [[#SIZE]] bytes (tgt:0x{{.*}}) -> (hst:0x{{0*}}[[#HOST_ADDR]])
// clang-format on
diff --git a/offload/test/mapping/map_ordering_tgt_exit_data_from_delete_assumedsize.c b/offload/test/mapping/map_ordering_tgt_exit_data_from_delete_assumedsize.c
index 963e709411c92..93284bdc5681a 100644
--- a/offload/test/mapping/map_ordering_tgt_exit_data_from_delete_assumedsize.c
+++ b/offload/test/mapping/map_ordering_tgt_exit_data_from_delete_assumedsize.c
@@ -32,7 +32,7 @@ int main() {
#pragma omp target exit data map(from : p2x[0]) map(delete : p1x[ : ])
// clang-format off
- // DEBUG: omptarget --> Pointer HstPtr=0x[[#%x,HOST_ADDR:]] falls within a range previously marked for deletion
+ // DEBUG: omptarget --> Pointer HstPtr=0x[[#%x,HOST_ADDR:]] falls within a range previously released
// DEBUG: omptarget --> Moving {{.*}} bytes (tgt:0x{{.*}}) -> (hst:0x{{0*}}[[#HOST_ADDR]])
// clang-format on
diff --git a/offload/test/mapping/map_ordering_tgt_exit_data_from_mapper_overlap.c b/offload/test/mapping/map_ordering_tgt_exit_data_from_mapper_overlap.c
new file mode 100644
index 0000000000000..af14920ab9898
--- /dev/null
+++ b/offload/test/mapping/map_ordering_tgt_exit_data_from_mapper_overlap.c
@@ -0,0 +1,49 @@
+// RUN: %libomptarget-compile-generic
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefix=CHECK
+// RUN: env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefix=DEBUG
+// REQUIRES: libomptarget-debug
+
+// The test ensures that the FROM transfer for the full "s1" is performed, and
+// not just the FROM done via the mapper of s1.s2.
+
+#include <stdio.h>
+
+typedef struct {
+ int a;
+ int b;
+} S2;
+
+#pragma omp declare mapper(my_mapper : S2 s2) map(tofrom : s2.a)
+
+typedef struct {
+ S2 s2;
+ int c;
+ int d;
+} S1;
+
+S1 s1;
+
+int main() {
+#pragma omp target enter data map(alloc : s1)
+
+#pragma omp target map(present, alloc : s1)
+ {
+ s1.s2.a = 111;
+ s1.s2.b = 222;
+ s1.c = 333;
+ s1.d = 444;
+ }
+
+ // clang-format off
+ // DEBUG: omptarget --> Tracking released entry: HstPtr=0x[[#%x,HOST_ADDR:]], Size=[[#%u,SIZE:]], ForceDelete=0
+ // DEBUG: omptarget --> Moving {{.*}} bytes (tgt:0x{{.*}}) -> (hst:0x{{.*}})
+ // DEBUG: omptarget --> Pointer HstPtr=0x{{0*}}[[#HOST_ADDR]] falls within a range previously released
+ // DEBUG: omptarget --> Moving [[#SIZE]] bytes (tgt:0x{{.*}}) -> (hst:0x{{0*}}[[#HOST_ADDR]])
+ // clang-format on
+#pragma omp target exit data map(from : s1) map(mapper(my_mapper), from : s1.s2)
+
+ // CHECK: 111 222 333 444
+ printf("%d %d %d %d\n", s1.s2.a, s1.s2.b, s1.c, s1.d);
+}
More information about the llvm-commits
mailing list