[Openmp-commits] [openmp] ec1ebcd - [OpenMP][OpenACC] Implement `ompx_hold` map type modifier extension in runtime (2/2)
Joel E. Denny via Openmp-commits
openmp-commits at lists.llvm.org
Tue Aug 31 13:15:10 PDT 2021
Author: Joel E. Denny
Date: 2021-08-31T16:13:49-04:00
New Revision: ec1ebcd30258329666dc89e6e745bb9de2e8fd13
URL: https://github.com/llvm/llvm-project/commit/ec1ebcd30258329666dc89e6e745bb9de2e8fd13
DIFF: https://github.com/llvm/llvm-project/commit/ec1ebcd30258329666dc89e6e745bb9de2e8fd13.diff
LOG: [OpenMP][OpenACC] Implement `ompx_hold` map type modifier extension in runtime (2/2)
This patch implements OpenMP runtime support for an original OpenMP
extension we have developed to support OpenACC: the `ompx_hold` map
type modifier. The previous patch in this series, D106509, implements
Clang support and documents the new functionality in detail.
Reviewed By: grokos
Differential Revision: https://reviews.llvm.org/D106510
Added:
openmp/libomptarget/test/mapping/ompx_hold/omp_target_disassociate_ptr.c
openmp/libomptarget/test/mapping/ompx_hold/struct.c
openmp/libomptarget/test/mapping/ompx_hold/target-data.c
openmp/libomptarget/test/mapping/ompx_hold/target.c
Modified:
openmp/libomptarget/include/omptarget.h
openmp/libomptarget/src/api.cpp
openmp/libomptarget/src/device.cpp
openmp/libomptarget/src/device.h
openmp/libomptarget/src/omptarget.cpp
openmp/libomptarget/src/private.h
openmp/libomptarget/test/offloading/info.c
Removed:
################################################################################
diff --git a/openmp/libomptarget/include/omptarget.h b/openmp/libomptarget/include/omptarget.h
index dfddb59fa91d9..c5b75d900d221 100644
--- a/openmp/libomptarget/include/omptarget.h
+++ b/openmp/libomptarget/include/omptarget.h
@@ -56,6 +56,10 @@ enum tgt_map_type {
OMP_TGT_MAPTYPE_CLOSE = 0x400,
// runtime error if not already allocated
OMP_TGT_MAPTYPE_PRESENT = 0x1000,
+ // use a separate reference counter so that the data cannot be unmapped within
+ // the structured region
+ // This is an OpenMP extension for the sake of OpenACC support.
+ OMP_TGT_MAPTYPE_OMPX_HOLD = 0x2000,
// descriptor for non-contiguous target-update
OMP_TGT_MAPTYPE_NON_CONTIG = 0x100000000000,
// member of struct, member given by [16 MSBs] - 1
diff --git a/openmp/libomptarget/src/api.cpp b/openmp/libomptarget/src/api.cpp
index ed641106c3a09..849ce3211ef7d 100644
--- a/openmp/libomptarget/src/api.cpp
+++ b/openmp/libomptarget/src/api.cpp
@@ -106,7 +106,8 @@ EXTERN int omp_target_is_present(const void *ptr, int device_num) {
bool IsLast; // not used
bool IsHostPtr;
void *TgtPtr = Device.getTgtPtrBegin(const_cast<void *>(ptr), 0, IsLast,
- false, IsHostPtr);
+ /*UpdateRefCount=*/false,
+ /*UseHoldRefCount=*/false, IsHostPtr);
int rc = (TgtPtr != NULL);
// Under unified memory the host pointer can be returned by the
// getTgtPtrBegin() function which means that there is no device
diff --git a/openmp/libomptarget/src/device.cpp b/openmp/libomptarget/src/device.cpp
index fd7c73df722c0..ff5b2882b46f6 100644
--- a/openmp/libomptarget/src/device.cpp
+++ b/openmp/libomptarget/src/device.cpp
@@ -82,14 +82,16 @@ int DeviceTy::associatePtr(void *HstPtrBegin, void *TgtPtrBegin, int64_t Size) {
/*HstPtrBase=*/(uintptr_t)HstPtrBegin,
/*HstPtrBegin=*/(uintptr_t)HstPtrBegin,
/*HstPtrEnd=*/(uintptr_t)HstPtrBegin + Size,
- /*TgtPtrBegin=*/(uintptr_t)TgtPtrBegin, /*Name=*/nullptr,
+ /*TgtPtrBegin=*/(uintptr_t)TgtPtrBegin,
+ /*UseHoldRefCount=*/false, /*Name=*/nullptr,
/*IsRefCountINF=*/true)
.first;
DP("Creating new map entry: HstBase=" DPxMOD ", HstBegin=" DPxMOD
- ", HstEnd=" DPxMOD ", TgtBegin=" DPxMOD ", RefCount=%s\n",
+ ", HstEnd=" DPxMOD ", TgtBegin=" DPxMOD ", DynRefCount=%s, "
+ "HoldRefCount=%s\n",
DPxPTR(newEntry.HstPtrBase), DPxPTR(newEntry.HstPtrBegin),
DPxPTR(newEntry.HstPtrEnd), DPxPTR(newEntry.TgtPtrBegin),
- newEntry.refCountToStr().c_str());
+ newEntry.dynRefCountToStr().c_str(), newEntry.holdRefCountToStr().c_str());
(void)newEntry;
DataMapMtx.unlock();
@@ -103,7 +105,13 @@ int DeviceTy::disassociatePtr(void *HstPtrBegin) {
auto search = HostDataToTargetMap.find(HstPtrBeginTy{(uintptr_t)HstPtrBegin});
if (search != HostDataToTargetMap.end()) {
// Mapping exists
- if (search->isRefCountInf()) {
+ if (search->getHoldRefCount()) {
+ // This is based on OpenACC 3.1, sec 3.2.33 "acc_unmap_data", L3656-3657:
+ // "It is an error to call acc_unmap_data if the structured reference
+ // count for the pointer is not zero."
+ REPORT("Trying to disassociate a pointer with a non-zero hold reference "
+ "count\n");
+ } else if (search->isDynRefCountInf()) {
DP("Association found, removing it\n");
HostDataToTargetMap.erase(search);
DataMapMtx.unlock();
@@ -112,11 +120,12 @@ int DeviceTy::disassociatePtr(void *HstPtrBegin) {
REPORT("Trying to disassociate a pointer which was not mapped via "
"omp_target_associate_ptr\n");
}
+ } else {
+ REPORT("Association not found\n");
}
// Mapping not found
DataMapMtx.unlock();
- REPORT("Association not found\n");
return OFFLOAD_FAIL;
}
@@ -171,7 +180,7 @@ DeviceTy::getTargetPointer(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
map_var_info_t HstPtrName, MoveDataStateTy MoveData,
bool IsImplicit, bool UpdateRefCount,
bool HasCloseModifier, bool HasPresentModifier,
- AsyncInfoTy &AsyncInfo) {
+ bool HasHoldModifier, AsyncInfoTy &AsyncInfo) {
void *TargetPointer = nullptr;
bool IsHostPtr = false;
bool IsNew = false;
@@ -188,21 +197,26 @@ DeviceTy::getTargetPointer(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
if (LR.Flags.IsContained ||
((LR.Flags.ExtendsBefore || LR.Flags.ExtendsAfter) && IsImplicit)) {
auto &HT = *LR.Entry;
- assert(HT.getRefCount() > 0 && "expected existing RefCount > 0");
- if (UpdateRefCount)
+ const char *RefCountAction;
+ assert(HT.getTotalRefCount() > 0 && "expected existing RefCount > 0");
+ if (UpdateRefCount) {
// After this, RefCount > 1.
- HT.incRefCount();
- else
+ HT.incRefCount(HasHoldModifier);
+ RefCountAction = " (incremented)";
+ } else {
// It might have been allocated with the parent, but it's still new.
- IsNew = HT.getRefCount() == 1;
+ IsNew = HT.getTotalRefCount() == 1;
+ RefCountAction = " (update suppressed)";
+ }
+ const char *DynRefCountAction = HasHoldModifier ? "" : RefCountAction;
+ const char *HoldRefCountAction = HasHoldModifier ? RefCountAction : "";
uintptr_t Ptr = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin);
INFO(OMP_INFOTYPE_MAPPING_EXISTS, DeviceID,
"Mapping exists%s with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD
- ", "
- "Size=%" PRId64 ", RefCount=%s (%s), Name=%s\n",
+ ", Size=%" PRId64 ", DynRefCount=%s%s, HoldRefCount=%s%s, Name=%s\n",
(IsImplicit ? " (implicit)" : ""), DPxPTR(HstPtrBegin), DPxPTR(Ptr),
- Size, HT.refCountToStr().c_str(),
- UpdateRefCount ? "incremented" : "update suppressed",
+ Size, HT.dynRefCountToStr().c_str(), DynRefCountAction,
+ HT.holdRefCountToStr().c_str(), HoldRefCountAction,
(HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "unknown");
TargetPointer = (void *)Ptr;
} else if ((LR.Flags.ExtendsBefore || LR.Flags.ExtendsAfter) && !IsImplicit) {
@@ -245,13 +259,15 @@ DeviceTy::getTargetPointer(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
uintptr_t Ptr = (uintptr_t)allocData(Size, HstPtrBegin);
Entry = HostDataToTargetMap
.emplace((uintptr_t)HstPtrBase, (uintptr_t)HstPtrBegin,
- (uintptr_t)HstPtrBegin + Size, Ptr, HstPtrName)
+ (uintptr_t)HstPtrBegin + Size, Ptr, HasHoldModifier,
+ HstPtrName)
.first;
INFO(OMP_INFOTYPE_MAPPING_CHANGED, DeviceID,
"Creating new map entry with "
"HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", Size=%ld, "
- "RefCount=%s, Name=%s\n",
- DPxPTR(HstPtrBegin), DPxPTR(Ptr), Size, Entry->refCountToStr().c_str(),
+ "DynRefCount=%s, HoldRefCount=%s, Name=%s\n",
+ DPxPTR(HstPtrBegin), DPxPTR(Ptr), Size,
+ Entry->dynRefCountToStr().c_str(), Entry->holdRefCountToStr().c_str(),
(HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "unknown");
TargetPointer = (void *)Ptr;
}
@@ -295,8 +311,9 @@ DeviceTy::getTargetPointer(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
// Return the target pointer begin (where the data will be moved).
// Decrement the reference counter if called from targetDataEnd.
void *DeviceTy::getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast,
- bool UpdateRefCount, bool &IsHostPtr,
- bool MustContain, bool ForceDelete) {
+ bool UpdateRefCount, bool UseHoldRefCount,
+ bool &IsHostPtr, bool MustContain,
+ bool ForceDelete) {
void *rc = NULL;
IsHostPtr = false;
IsLast = false;
@@ -306,35 +323,39 @@ void *DeviceTy::getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast,
if (lr.Flags.IsContained ||
(!MustContain && (lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter))) {
auto &HT = *lr.Entry;
- // We do not decrement the reference count to zero here. deallocTgtPtr does
- // that atomically with removing the mapping. Otherwise, before this thread
+ // We do not zero the total reference count here. deallocTgtPtr does that
+ // atomically with removing the mapping. Otherwise, before this thread
// removed the mapping in deallocTgtPtr, another thread could retrieve the
// mapping, increment and decrement back to zero, and then both threads
// would try to remove the mapping, resulting in a double free.
- IsLast = HT.decShouldRemove(ForceDelete);
+ IsLast = HT.decShouldRemove(UseHoldRefCount, ForceDelete);
const char *RefCountAction;
if (!UpdateRefCount) {
- RefCountAction = "update suppressed";
+ RefCountAction = " (update suppressed)";
} else if (ForceDelete) {
- HT.resetRefCount();
- assert(IsLast == HT.decShouldRemove() &&
+ HT.resetRefCount(UseHoldRefCount);
+ assert(IsLast == HT.decShouldRemove(UseHoldRefCount) &&
"expected correct IsLast prediction for reset");
if (IsLast)
- RefCountAction = "reset, deferred final decrement";
- else
- RefCountAction = "reset";
+ RefCountAction = " (reset, deferred final decrement)";
+ else {
+ HT.decRefCount(UseHoldRefCount);
+ RefCountAction = " (reset)";
+ }
} else if (IsLast) {
- RefCountAction = "deferred final decrement";
+ RefCountAction = " (deferred final decrement)";
} else {
- RefCountAction = "decremented";
- HT.decRefCount();
+ HT.decRefCount(UseHoldRefCount);
+ RefCountAction = " (decremented)";
}
+ const char *DynRefCountAction = UseHoldRefCount ? "" : RefCountAction;
+ const char *HoldRefCountAction = UseHoldRefCount ? RefCountAction : "";
uintptr_t tp = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin);
INFO(OMP_INFOTYPE_MAPPING_EXISTS, DeviceID,
"Mapping exists with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", "
- "Size=%" PRId64 ", RefCount=%s (%s)\n",
- DPxPTR(HstPtrBegin), DPxPTR(tp), Size, HT.refCountToStr().c_str(),
- RefCountAction);
+ "Size=%" PRId64 ", DynRefCount=%s%s, HoldRefCount=%s%s\n",
+ DPxPTR(HstPtrBegin), DPxPTR(tp), Size, HT.dynRefCountToStr().c_str(),
+ DynRefCountAction, HT.holdRefCountToStr().c_str(), HoldRefCountAction);
rc = (void *)tp;
} else if (PM->RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) {
// If the value isn't found in the mapping and unified shared memory
@@ -366,7 +387,7 @@ void *DeviceTy::getTgtPtrBegin(void *HstPtrBegin, int64_t Size) {
}
int DeviceTy::deallocTgtPtr(void *HstPtrBegin, int64_t Size,
- bool HasCloseModifier) {
+ bool HasCloseModifier, bool HasHoldModifier) {
if (PM->RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
!HasCloseModifier)
return OFFLOAD_SUCCESS;
@@ -376,7 +397,7 @@ int DeviceTy::deallocTgtPtr(void *HstPtrBegin, int64_t Size,
LookupResult lr = lookupMapping(HstPtrBegin, Size);
if (lr.Flags.IsContained || lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) {
auto &HT = *lr.Entry;
- if (HT.decRefCount() == 0) {
+ if (HT.decRefCount(HasHoldModifier) == 0) {
DP("Deleting tgt data " DPxMOD " of size %" PRId64 "\n",
DPxPTR(HT.TgtPtrBegin), Size);
deleteData((void *)HT.TgtPtrBegin);
diff --git a/openmp/libomptarget/src/device.h b/openmp/libomptarget/src/device.h
index 58c6316ff6c32..78a8e274857a8 100644
--- a/openmp/libomptarget/src/device.h
+++ b/openmp/libomptarget/src/device.h
@@ -50,9 +50,30 @@ struct HostDataToTargetTy {
uintptr_t TgtPtrBegin; // target info.
private:
- /// use mutable to allow modification via std::set iterator which is const.
- mutable uint64_t RefCount;
+ /// The dynamic reference count is the standard reference count as of OpenMP
+ /// 4.5. The hold reference count is an OpenMP extension for the sake of
+ /// OpenACC support.
+ ///
+ /// The 'ompx_hold' map type modifier is permitted only on "omp target" and
+ /// "omp target data", and "delete" is permitted only on "omp target exit
+ /// data" and associated runtime library routines. As a result, we really
+ /// need to implement "reset" functionality only for the dynamic reference
+ /// counter. Likewise, only the dynamic reference count can be infinite
+ /// because, for example, omp_target_associate_ptr and "omp declare target
+ /// link" operate only on it. Nevertheless, it's actually easier to follow
+ /// the code (and requires less assertions for special cases) when we just
+ /// implement these features generally across both reference counters here.
+ /// Thus, it's the users of this class that impose those restrictions.
+ ///
+ /// Use mutable to allow modification via std::set iterator which is const.
+ ///@{
+ mutable uint64_t DynRefCount;
+ mutable uint64_t HoldRefCount;
+ ///@}
static const uint64_t INFRefCount = ~(uint64_t)0;
+ static std::string refCountToStr(uint64_t RefCount) {
+ return RefCount == INFRefCount ? "INF" : std::to_string(RefCount);
+ }
/// This mutex will be locked when data movement is issued. For targets that
/// doesn't support async data movement, this mutex can guarantee that after
/// it is released, memory region on the target is update to date. For targets
@@ -63,50 +84,82 @@ struct HostDataToTargetTy {
public:
HostDataToTargetTy(uintptr_t BP, uintptr_t B, uintptr_t E, uintptr_t TB,
- map_var_info_t Name = nullptr, bool IsINF = false)
+ bool UseHoldRefCount, map_var_info_t Name = nullptr,
+ bool IsINF = false)
: HstPtrBase(BP), HstPtrBegin(B), HstPtrEnd(E), HstPtrName(Name),
- TgtPtrBegin(TB), RefCount(IsINF ? INFRefCount : 1),
+ TgtPtrBegin(TB), DynRefCount(UseHoldRefCount ? 0
+ : IsINF ? INFRefCount
+ : 1),
+ HoldRefCount(!UseHoldRefCount ? 0
+ : IsINF ? INFRefCount
+ : 1),
UpdateMtx(std::make_shared<std::mutex>()) {}
- uint64_t getRefCount() const { return RefCount; }
+ /// Get the total reference count. This is smarter than just getDynRefCount()
+ /// + getHoldRefCount() because it handles the case where at least one is
+ /// infinity and the other is non-zero.
+ uint64_t getTotalRefCount() const {
+ if (DynRefCount == INFRefCount || HoldRefCount == INFRefCount)
+ return INFRefCount;
+ return DynRefCount + HoldRefCount;
+ }
+
+ /// Get the dynamic reference count.
+ uint64_t getDynRefCount() const { return DynRefCount; }
- uint64_t resetRefCount() const {
- if (RefCount != INFRefCount)
- RefCount = 1;
+ /// Get the hold reference count.
+ uint64_t getHoldRefCount() const { return HoldRefCount; }
- return RefCount;
+ /// Reset the specified reference count unless it's infinity. Reset to 1
+ /// (even if currently 0) so it can be followed by a decrement.
+ void resetRefCount(bool UseHoldRefCount) const {
+ uint64_t &ThisRefCount = UseHoldRefCount ? HoldRefCount : DynRefCount;
+ if (ThisRefCount != INFRefCount)
+ ThisRefCount = 1;
}
- uint64_t incRefCount() const {
- if (RefCount != INFRefCount) {
- ++RefCount;
- assert(RefCount < INFRefCount && "refcount overflow");
+ /// Increment the specified reference count unless it's infinity.
+ void incRefCount(bool UseHoldRefCount) const {
+ uint64_t &ThisRefCount = UseHoldRefCount ? HoldRefCount : DynRefCount;
+ if (ThisRefCount != INFRefCount) {
+ ++ThisRefCount;
+ assert(ThisRefCount < INFRefCount && "refcount overflow");
}
-
- return RefCount;
}
- uint64_t decRefCount() const {
- if (RefCount != INFRefCount) {
- assert(RefCount > 0 && "refcount underflow");
- --RefCount;
+ /// Decrement the specified reference count unless it's infinity or zero, and
+ /// return the total reference count.
+ uint64_t decRefCount(bool UseHoldRefCount) const {
+ uint64_t &ThisRefCount = UseHoldRefCount ? HoldRefCount : DynRefCount;
+ uint64_t OtherRefCount = UseHoldRefCount ? DynRefCount : HoldRefCount;
+ if (ThisRefCount != INFRefCount) {
+ if (ThisRefCount > 0)
+ --ThisRefCount;
+ else
+ assert(OtherRefCount > 0 && "total refcount underflow");
}
-
- return RefCount;
+ return getTotalRefCount();
}
- bool isRefCountInf() const { return RefCount == INFRefCount; }
+ /// Is the dynamic (and thus the total) reference count infinite?
+ bool isDynRefCountInf() const { return DynRefCount == INFRefCount; }
- std::string refCountToStr() const {
- return isRefCountInf() ? "INF" : std::to_string(getRefCount());
- }
+ /// Convert the dynamic reference count to a debug string.
+ std::string dynRefCountToStr() const { return refCountToStr(DynRefCount); }
+
+ /// Convert the hold reference count to a debug string.
+ std::string holdRefCountToStr() const { return refCountToStr(HoldRefCount); }
- /// Should one decrement of the reference count (after resetting it if
- /// \c AfterReset) remove this mapping?
- bool decShouldRemove(bool AfterReset = false) const {
+ /// Should one decrement of the specified reference count (after resetting it
+ /// if \c AfterReset) remove this mapping?
+ bool decShouldRemove(bool UseHoldRefCount, bool AfterReset = false) const {
+ uint64_t ThisRefCount = UseHoldRefCount ? HoldRefCount : DynRefCount;
+ uint64_t OtherRefCount = UseHoldRefCount ? DynRefCount : HoldRefCount;
+ if (OtherRefCount > 0)
+ return false;
if (AfterReset)
- return !isRefCountInf();
- return getRefCount() == 1;
+ return ThisRefCount != INFRefCount;
+ return ThisRefCount == 1;
}
void lock() const { UpdateMtx->lock(); }
@@ -223,13 +276,15 @@ struct DeviceTy {
getTargetPointer(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
map_var_info_t HstPtrName, MoveDataStateTy MoveData,
bool IsImplicit, bool UpdateRefCount, bool HasCloseModifier,
- bool HasPresentModifier, AsyncInfoTy &AsyncInfo);
+ bool HasPresentModifier, bool HasHoldModifier,
+ AsyncInfoTy &AsyncInfo);
void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size);
void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast,
- bool UpdateRefCount, bool &IsHostPtr,
- bool MustContain = false, bool ForceDelete = false);
- int deallocTgtPtr(void *TgtPtrBegin, int64_t Size,
- bool HasCloseModifier = false);
+ bool UpdateRefCount, bool UseHoldRefCount,
+ bool &IsHostPtr, bool MustContain = false,
+ bool ForceDelete = false);
+ int deallocTgtPtr(void *TgtPtrBegin, int64_t Size, bool HasCloseModifier,
+ bool HasHoldModifier);
int associatePtr(void *HstPtrBegin, void *TgtPtrBegin, int64_t Size);
int disassociatePtr(void *HstPtrBegin);
diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp
index 5f084e75ce84c..086503a6be6d1 100644
--- a/openmp/libomptarget/src/omptarget.cpp
+++ b/openmp/libomptarget/src/omptarget.cpp
@@ -157,7 +157,8 @@ static int InitLibrary(DeviceTy &Device) {
(uintptr_t)CurrHostEntry->addr /*HstPtrBase*/,
(uintptr_t)CurrHostEntry->addr /*HstPtrBegin*/,
(uintptr_t)CurrHostEntry->addr + CurrHostEntry->size /*HstPtrEnd*/,
- (uintptr_t)CurrDeviceEntry->addr /*TgtPtrBegin*/, nullptr,
+ (uintptr_t)CurrDeviceEntry->addr /*TgtPtrBegin*/,
+ false /*UseHoldRefCount*/, nullptr /*Name*/,
true /*IsRefCountINF*/);
}
}
@@ -465,6 +466,7 @@ int targetDataBegin(ident_t *loc, DeviceTy &Device, int32_t arg_num,
// a close map modifier was associated with a map that contained a to.
bool HasCloseModifier = arg_types[i] & OMP_TGT_MAPTYPE_CLOSE;
bool HasPresentModifier = arg_types[i] & OMP_TGT_MAPTYPE_PRESENT;
+ bool HasHoldModifier = arg_types[i] & OMP_TGT_MAPTYPE_OMPX_HOLD;
// UpdateRef is based on MEMBER_OF instead of TARGET_PARAM because if we
// have reached this point via __tgt_target_data_begin and not __tgt_target
// then no argument is marked as TARGET_PARAM ("omp target data map" is not
@@ -490,7 +492,7 @@ int targetDataBegin(ident_t *loc, DeviceTy &Device, int32_t arg_num,
Pointer_TPR = Device.getTargetPointer(
HstPtrBase, HstPtrBase, sizeof(void *), nullptr,
MoveDataStateTy::NONE, IsImplicit, UpdateRef, HasCloseModifier,
- HasPresentModifier, AsyncInfo);
+ HasPresentModifier, HasHoldModifier, AsyncInfo);
PointerTgtPtrBegin = Pointer_TPR.TargetPointer;
IsHostPtr = Pointer_TPR.Flags.IsHostPointer;
if (!PointerTgtPtrBegin) {
@@ -522,7 +524,8 @@ int targetDataBegin(ident_t *loc, DeviceTy &Device, int32_t arg_num,
auto TPR = Device.getTargetPointer(
HstPtrBegin, HstPtrBase, data_size, HstPtrName, MoveData, IsImplicit,
- UpdateRef, HasCloseModifier, HasPresentModifier, AsyncInfo);
+ UpdateRef, HasCloseModifier, HasPresentModifier, HasHoldModifier,
+ AsyncInfo);
void *TgtPtrBegin = TPR.TargetPointer;
IsHostPtr = TPR.Flags.IsHostPointer;
// If data_size==0, then the argument could be a zero-length pointer to
@@ -608,10 +611,13 @@ struct DeallocTgtPtrInfo {
int64_t DataSize;
/// Whether it has \p close modifier
bool HasCloseModifier;
+ /// Whether it has \p ompx_hold modifier
+ bool HasHoldModifier;
- DeallocTgtPtrInfo(void *HstPtr, int64_t Size, bool HasCloseModifier)
- : HstPtrBegin(HstPtr), DataSize(Size),
- HasCloseModifier(HasCloseModifier) {}
+ DeallocTgtPtrInfo(void *HstPtr, int64_t Size, bool HasCloseModifier,
+ bool HasHoldModifier)
+ : HstPtrBegin(HstPtr), DataSize(Size), HasCloseModifier(HasCloseModifier),
+ HasHoldModifier(HasHoldModifier) {}
};
} // namespace
@@ -678,11 +684,12 @@ int targetDataEnd(ident_t *loc, DeviceTy &Device, int32_t ArgNum,
bool ForceDelete = ArgTypes[I] & OMP_TGT_MAPTYPE_DELETE;
bool HasCloseModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_CLOSE;
bool HasPresentModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_PRESENT;
+ bool HasHoldModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_OMPX_HOLD;
// If PTR_AND_OBJ, HstPtrBegin is address of pointee
- void *TgtPtrBegin =
- Device.getTgtPtrBegin(HstPtrBegin, DataSize, IsLast, UpdateRef,
- IsHostPtr, !IsImplicit, ForceDelete);
+ void *TgtPtrBegin = Device.getTgtPtrBegin(
+ HstPtrBegin, DataSize, IsLast, UpdateRef, HasHoldModifier, IsHostPtr,
+ !IsImplicit, ForceDelete);
if (!TgtPtrBegin && (DataSize || HasPresentModifier)) {
DP("Mapping does not exist (%s)\n",
(HasPresentModifier ? "'present' map type modifier" : "ignored"));
@@ -799,7 +806,8 @@ int targetDataEnd(ident_t *loc, DeviceTy &Device, int32_t ArgNum,
// Add pointer to the buffer for later deallocation
if (DelEntry)
- DeallocTgtPtrs.emplace_back(HstPtrBegin, DataSize, HasCloseModifier);
+ DeallocTgtPtrs.emplace_back(HstPtrBegin, DataSize, HasCloseModifier,
+ HasHoldModifier);
}
}
@@ -816,7 +824,7 @@ int targetDataEnd(ident_t *loc, DeviceTy &Device, int32_t ArgNum,
if (FromMapperBase && FromMapperBase == Info.HstPtrBegin)
continue;
Ret = Device.deallocTgtPtr(Info.HstPtrBegin, Info.DataSize,
- Info.HasCloseModifier);
+ Info.HasCloseModifier, Info.HasHoldModifier);
if (Ret != OFFLOAD_SUCCESS) {
REPORT("Deallocating data from device failed.\n");
return OFFLOAD_FAIL;
@@ -831,8 +839,9 @@ static int targetDataContiguous(ident_t *loc, DeviceTy &Device, void *ArgsBase,
int64_t ArgType, AsyncInfoTy &AsyncInfo) {
TIMESCOPE_WITH_IDENT(loc);
bool IsLast, IsHostPtr;
- void *TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, ArgSize, IsLast, false,
- IsHostPtr, /*MustContain=*/true);
+ void *TgtPtrBegin = Device.getTgtPtrBegin(
+ HstPtrBegin, ArgSize, IsLast, /*UpdateRefCount=*/false,
+ /*UseHoldRefCount=*/false, IsHostPtr, /*MustContain=*/true);
if (!TgtPtrBegin) {
DP("hst data:" DPxMOD " not found, becomes a noop\n", DPxPTR(HstPtrBegin));
if (ArgType & OMP_TGT_MAPTYPE_PRESENT) {
@@ -1291,8 +1300,9 @@ static int processDataBefore(ident_t *loc, int64_t DeviceId, void *HostPtr,
uint64_t Delta = (uint64_t)HstPtrBegin - (uint64_t)HstPtrBase;
void *TgtPtrBegin = (void *)((uintptr_t)TgtPtrBase + Delta);
void *&PointerTgtPtrBegin = AsyncInfo.getVoidPtrLocation();
- PointerTgtPtrBegin = Device.getTgtPtrBegin(HstPtrVal, ArgSizes[I],
- IsLast, false, IsHostPtr);
+ PointerTgtPtrBegin = Device.getTgtPtrBegin(
+ HstPtrVal, ArgSizes[I], IsLast, /*UpdateRefCount=*/false,
+ /*UseHoldRefCount=*/false, IsHostPtr);
if (!PointerTgtPtrBegin) {
DP("No lambda captured variable mapped (" DPxMOD ") - ignored\n",
DPxPTR(HstPtrVal));
@@ -1348,7 +1358,8 @@ static int processDataBefore(ident_t *loc, int64_t DeviceId, void *HostPtr,
if (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)
HstPtrBase = *reinterpret_cast<void **>(HstPtrBase);
TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, ArgSizes[I], IsLast,
- false, IsHostPtr);
+ /*UpdateRefCount=*/false,
+ /*UseHoldRefCount=*/false, IsHostPtr);
TgtBaseOffset = (intptr_t)HstPtrBase - (intptr_t)HstPtrBegin;
#ifdef OMPTARGET_DEBUG
void *TgtPtrBase = (void *)((intptr_t)TgtPtrBegin + TgtBaseOffset);
diff --git a/openmp/libomptarget/src/private.h b/openmp/libomptarget/src/private.h
index 05996eb5cdf39..bbccf9e5c475a 100644
--- a/openmp/libomptarget/src/private.h
+++ b/openmp/libomptarget/src/private.h
@@ -111,16 +111,18 @@ static inline void dumpTargetPointerMappings(const ident_t *Loc,
INFO(OMP_INFOTYPE_ALL, Device.DeviceID,
"OpenMP Host-Device pointer mappings after block at %s:%d:%d:\n",
Kernel.getFilename(), Kernel.getLine(), Kernel.getColumn());
- INFO(OMP_INFOTYPE_ALL, Device.DeviceID, "%-18s %-18s %s %s %s\n", "Host Ptr",
- "Target Ptr", "Size (B)", "RefCount", "Declaration");
+ INFO(OMP_INFOTYPE_ALL, Device.DeviceID, "%-18s %-18s %s %s %s %s\n",
+ "Host Ptr", "Target Ptr", "Size (B)", "DynRefCount", "HoldRefCount",
+ "Declaration");
Device.DataMapMtx.lock();
for (const auto &HostTargetMap : Device.HostDataToTargetMap) {
SourceInfo Info(HostTargetMap.HstPtrName);
INFO(OMP_INFOTYPE_ALL, Device.DeviceID,
- DPxMOD " " DPxMOD " %-8" PRIuPTR " %-8s %s at %s:%d:%d\n",
+ DPxMOD " " DPxMOD " %-8" PRIuPTR " %-11s %-12s %s at %s:%d:%d\n",
DPxPTR(HostTargetMap.HstPtrBegin), DPxPTR(HostTargetMap.TgtPtrBegin),
HostTargetMap.HstPtrEnd - HostTargetMap.HstPtrBegin,
- HostTargetMap.refCountToStr().c_str(), Info.getName(),
+ HostTargetMap.dynRefCountToStr().c_str(),
+ HostTargetMap.holdRefCountToStr().c_str(), Info.getName(),
Info.getFilename(), Info.getLine(), Info.getColumn());
}
Device.DataMapMtx.unlock();
diff --git a/openmp/libomptarget/test/mapping/ompx_hold/omp_target_disassociate_ptr.c b/openmp/libomptarget/test/mapping/ompx_hold/omp_target_disassociate_ptr.c
new file mode 100644
index 0000000000000..6bd0e91991772
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/ompx_hold/omp_target_disassociate_ptr.c
@@ -0,0 +1,68 @@
+// omp_target_disassociate_ptr should always fail if the hold reference count is
+// non-zero, regardless of the dynamic reference count. When the latter is
+// finite, the implementation happens to choose to report the hold diagnostic.
+
+// RUN: %libomptarget-compile-generic -fopenmp-extensions
+// RUN: %not %libomptarget-run-generic 0 2>&1 | %fcheck-generic
+// RUN: %not %libomptarget-run-generic 1 2>&1 | %fcheck-generic
+// RUN: %not %libomptarget-run-generic inf 2>&1 | %fcheck-generic
+
+// RUN: %libomptarget-compile-generic -fopenmp-extensions -DHOLD_MORE
+// RUN: %not %libomptarget-run-generic 0 2>&1 | %fcheck-generic
+// RUN: %not %libomptarget-run-generic 1 2>&1 | %fcheck-generic
+// RUN: %not %libomptarget-run-generic inf 2>&1 | %fcheck-generic
+
+#include <omp.h>
+#include <stdio.h>
+#include <limits.h>
+#include <string.h>
+
+int main(int argc, char *argv[]) {
+ // Parse command line.
+ int DynRef;
+ if (argc != 2) {
+ fprintf(stderr, "bad arguments\n");
+ return 1;
+ }
+ if (0 == strcmp(argv[1], "inf"))
+ DynRef = INT_MAX;
+ else
+ DynRef = atoi(argv[1]);
+
+ // Allocate and set dynamic reference count as specified.
+ int DevNum = omp_get_default_device();
+ int X;
+ void *XDev = omp_target_alloc(sizeof X, DevNum);
+ if (!XDev) {
+ fprintf(stderr, "omp_target_alloc failed\n");
+ return 1;
+ }
+ if (DynRef == INT_MAX) {
+ if (omp_target_associate_ptr(&X, &XDev, sizeof X, 0, DevNum)) {
+ fprintf(stderr, "omp_target_associate_ptr failed\n");
+ return 1;
+ }
+ } else {
+ for (int I = 0; I < DynRef; ++I) {
+ #pragma omp target enter data map(alloc: X)
+ }
+ }
+
+ // Disassociate while hold reference count > 0.
+ int Status = 0;
+ #pragma omp target data map(ompx_hold,alloc: X)
+#if HOLD_MORE
+ #pragma omp target data map(ompx_hold,alloc: X)
+ #pragma omp target data map(ompx_hold,alloc: X)
+#endif
+ {
+ // CHECK: Libomptarget error: Trying to disassociate a pointer with a
+ // CHECK-SAME: non-zero hold reference count
+ // CHECK-NEXT: omp_target_disassociate_ptr failed
+ if (omp_target_disassociate_ptr(&X, DevNum)) {
+ fprintf(stderr, "omp_target_disassociate_ptr failed\n");
+ Status = 1;
+ }
+ }
+ return Status;
+}
diff --git a/openmp/libomptarget/test/mapping/ompx_hold/struct.c b/openmp/libomptarget/test/mapping/ompx_hold/struct.c
new file mode 100644
index 0000000000000..63c6ef1e68a29
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/ompx_hold/struct.c
@@ -0,0 +1,202 @@
+// RUN: %libomptarget-compile-generic -fopenmp-extensions
+// RUN: %libomptarget-run-generic | %fcheck-generic -strict-whitespace
+
+#include <omp.h>
+#include <stdio.h>
+
+#define CHECK_PRESENCE(Var1, Var2, Var3) \
+ printf(" presence of %s, %s, %s: %d, %d, %d\n", \
+ #Var1, #Var2, #Var3, \
+ omp_target_is_present(&(Var1), omp_get_default_device()), \
+ omp_target_is_present(&(Var2), omp_get_default_device()), \
+ omp_target_is_present(&(Var3), omp_get_default_device()))
+
+#define CHECK_VALUES(Var1, Var2) \
+ printf(" values of %s, %s: %d, %d\n", \
+ #Var1, #Var2, (Var1), (Var2))
+
+int main() {
+ struct S { int i; int j; } s;
+ // CHECK: presence of s, s.i, s.j: 0, 0, 0
+ CHECK_PRESENCE(s, s.i, s.j);
+
+ // =======================================================================
+ // Check that ompx_hold keeps entire struct present.
+
+ // -----------------------------------------------------------------------
+ // CHECK-LABEL: check:{{.*}}
+ printf("check: ompx_hold only on first member\n");
+ s.i = 20;
+ s.j = 30;
+ #pragma omp target data map(tofrom: s) map(ompx_hold,tofrom: s.i) \
+ map(tofrom: s.j)
+ {
+ // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
+ CHECK_PRESENCE(s, s.i, s.j);
+ #pragma omp target map(tofrom: s)
+ {
+ s.i = 21;
+ s.j = 31;
+ }
+ #pragma omp target exit data map(delete: s, s.i)
+ // ompx_hold on s.i applies to all of s.
+ // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
+ // CHECK-NEXT: values of s.i, s.j: 20, 30
+ CHECK_PRESENCE(s, s.i, s.j);
+ CHECK_VALUES(s.i, s.j);
+ }
+ // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
+ // CHECK-NEXT: values of s.i, s.j: 21, 31
+ CHECK_PRESENCE(s, s.i, s.j);
+ CHECK_VALUES(s.i, s.j);
+
+ // -----------------------------------------------------------------------
+ // CHECK-LABEL: check:{{.*}}
+ printf("check: ompx_hold only on last member\n");
+ s.i = 20;
+ s.j = 30;
+ #pragma omp target data map(tofrom: s) map(tofrom: s.i) \
+ map(ompx_hold,tofrom: s.j)
+ {
+ // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
+ CHECK_PRESENCE(s, s.i, s.j);
+ #pragma omp target map(tofrom: s)
+ {
+ s.i = 21;
+ s.j = 31;
+ }
+ #pragma omp target exit data map(delete: s, s.i)
+ // ompx_hold on s.j applies to all of s.
+ // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
+ // CHECK-NEXT: values of s.i, s.j: 20, 30
+ CHECK_PRESENCE(s, s.i, s.j);
+ CHECK_VALUES(s.i, s.j);
+ }
+ // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
+ // CHECK-NEXT: values of s.i, s.j: 21, 31
+ CHECK_PRESENCE(s, s.i, s.j);
+ CHECK_VALUES(s.i, s.j);
+
+ // -----------------------------------------------------------------------
+ // CHECK-LABEL: check:{{.*}}
+ printf("check: ompx_hold only on struct\n");
+ s.i = 20;
+ s.j = 30;
+ #pragma omp target data map(ompx_hold,tofrom: s) map(tofrom: s.i) \
+ map(tofrom: s.j)
+ {
+ // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
+ CHECK_PRESENCE(s, s.i, s.j);
+ #pragma omp target map(tofrom: s)
+ {
+ s.i = 21;
+ s.j = 31;
+ }
+ #pragma omp target exit data map(delete: s, s.i)
+ // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
+ // CHECK-NEXT: values of s.i, s.j: 20, 30
+ CHECK_PRESENCE(s, s.i, s.j);
+ CHECK_VALUES(s.i, s.j);
+ }
+ // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
+ // CHECK-NEXT: values of s.i, s.j: 21, 31
+ CHECK_PRESENCE(s, s.i, s.j);
+ CHECK_VALUES(s.i, s.j);
+
+ // =======================================================================
+ // Check that transfer to/from host checks reference count correctly.
+
+ // -----------------------------------------------------------------------
+ // CHECK-LABEL: check:{{.*}}
+ printf("check: parent DynRefCount=1 is not sufficient for transfer\n");
+ s.i = 20;
+ s.j = 30;
+ #pragma omp target data map(ompx_hold, tofrom: s)
+ #pragma omp target data map(ompx_hold, tofrom: s)
+ {
+ // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
+ CHECK_PRESENCE(s, s.i, s.j);
+ #pragma omp target map(from: s.i, s.j)
+ {
+ s.i = 21;
+ s.j = 31;
+ } // No transfer here even though parent's DynRefCount=1.
+ // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
+ // CHECK-NEXT: values of s.i, s.j: 20, 30
+ CHECK_PRESENCE(s, s.i, s.j);
+ CHECK_VALUES(s.i, s.j);
+ #pragma omp target map(to: s.i, s.j)
+ { // No transfer here even though parent's DynRefCount=1.
+ // CHECK-NEXT: values of s.i, s.j: 21, 31
+ CHECK_VALUES(s.i, s.j);
+ }
+ }
+ // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
+ // CHECK-NEXT: values of s.i, s.j: 21, 31
+ CHECK_PRESENCE(s, s.i, s.j);
+ CHECK_VALUES(s.i, s.j);
+
+ // -----------------------------------------------------------------------
+ // CHECK-LABEL: check:{{.*}}
+ printf("check: parent HoldRefCount=1 is not sufficient for transfer\n");
+ s.i = 20;
+ s.j = 30;
+ #pragma omp target data map(tofrom: s)
+ #pragma omp target data map(tofrom: s)
+ {
+ // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
+ CHECK_PRESENCE(s, s.i, s.j);
+ #pragma omp target map(ompx_hold, from: s.i, s.j)
+ {
+ s.i = 21;
+ s.j = 31;
+ } // No transfer here even though parent's HoldRefCount=1.
+ // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
+ // CHECK-NEXT: values of s.i, s.j: 20, 30
+ CHECK_PRESENCE(s, s.i, s.j);
+ CHECK_VALUES(s.i, s.j);
+ #pragma omp target map(ompx_hold, to: s.i, s.j)
+ { // No transfer here even though parent's HoldRefCount=1.
+ // CHECK-NEXT: values of s.i, s.j: 21, 31
+ CHECK_VALUES(s.i, s.j);
+ }
+ }
+ // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
+ // CHECK-NEXT: values of s.i, s.j: 21, 31
+ CHECK_PRESENCE(s, s.i, s.j);
+ CHECK_VALUES(s.i, s.j);
+
+ // -----------------------------------------------------------------------
+ // CHECK-LABEL: check:{{.*}}
+ //
+ // At the beginning of a region, if the parent's TotalRefCount=1, then the
+ // transfer should happen.
+ //
+ // At the end of a region, it also must be true that the reference count being
+ // decremented is the reference count that is 1.
+ printf("check: parent TotalRefCount=1 is not sufficient for transfer\n");
+ s.i = 20;
+ s.j = 30;
+ #pragma omp target data map(ompx_hold, tofrom: s)
+ {
+ // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
+ CHECK_PRESENCE(s, s.i, s.j);
+ #pragma omp target map(ompx_hold, tofrom: s.i, s.j)
+ {
+ s.i = 21;
+ s.j = 31;
+ }
+ #pragma omp target exit data map(from: s.i, s.j)
+ // No transfer here even though parent's TotalRefCount=1.
+ // CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
+ // CHECK-NEXT: values of s.i, s.j: 20, 30
+ CHECK_PRESENCE(s, s.i, s.j);
+ CHECK_VALUES(s.i, s.j);
+ }
+ // CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
+ // CHECK-NEXT: values of s.i, s.j: 21, 31
+ CHECK_PRESENCE(s, s.i, s.j);
+ CHECK_VALUES(s.i, s.j);
+
+ return 0;
+}
diff --git a/openmp/libomptarget/test/mapping/ompx_hold/target-data.c b/openmp/libomptarget/test/mapping/ompx_hold/target-data.c
new file mode 100644
index 0000000000000..154eb0358391d
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/ompx_hold/target-data.c
@@ -0,0 +1,236 @@
+// RUN: %libomptarget-compile-generic -fopenmp-extensions
+// RUN: %libomptarget-run-generic | %fcheck-generic -strict-whitespace
+
+#include <omp.h>
+#include <stdio.h>
+
+#define CHECK_PRESENCE(Var1, Var2, Var3) \
+ printf(" presence of %s, %s, %s: %d, %d, %d\n", \
+ #Var1, #Var2, #Var3, \
+ omp_target_is_present(&Var1, omp_get_default_device()), \
+ omp_target_is_present(&Var2, omp_get_default_device()), \
+ omp_target_is_present(&Var3, omp_get_default_device()))
+
+int main() {
+ int m, r, d;
+ // CHECK: presence of m, r, d: 0, 0, 0
+ CHECK_PRESENCE(m, r, d);
+
+ // -----------------------------------------------------------------------
+ // CHECK-NEXT: check:{{.*}}
+ printf("check: dyn>0, hold=0, dec/reset dyn=0\n");
+
+ // CHECK-NEXT: structured{{.*}}
+ printf(" structured dec of dyn\n");
+ #pragma omp target data map(tofrom: m) map(alloc: r, d)
+ {
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target data map(tofrom: m) map(alloc: r, d)
+ {
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ }
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ }
+ // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+ CHECK_PRESENCE(m, r, d);
+
+ // CHECK-NEXT: dynamic{{.*}}
+ printf(" dynamic dec/reset of dyn\n");
+ #pragma omp target data map(tofrom: m) map(alloc: r, d)
+ {
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target data map(tofrom: m) map(alloc: r, d)
+ {
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target exit data map(from: m) map(release: r)
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target exit data map(from: m) map(release: r) map(delete: d)
+ // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+ CHECK_PRESENCE(m, r, d);
+ }
+ // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target exit data map(from: m) map(release: r) map(delete: d)
+ // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+ CHECK_PRESENCE(m, r, d);
+ }
+ // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+ CHECK_PRESENCE(m, r, d);
+
+ // -----------------------------------------------------------------------
+ // CHECK: check:{{.*}}
+ printf("check: dyn=0, hold>0, dec/reset dyn=0, dec hold=0\n");
+
+ // Structured dec of dyn would require dyn>0.
+
+ // CHECK-NEXT: dynamic{{.*}}
+ printf(" dynamic dec/reset of dyn\n");
+ #pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+ {
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target data map(ompx_hold, tofrom: m) \
+ map(ompx_hold, alloc: r, d)
+ {
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target exit data map(from: m) map(release: r)
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target exit data map(from: m) map(release: r) map(delete: d)
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ }
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target exit data map(from: m) map(release: r) map(delete: d)
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ }
+ // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+ CHECK_PRESENCE(m, r, d);
+
+ // -----------------------------------------------------------------------
+ // CHECK: check:{{.*}}
+ printf("check: dyn>0, hold>0, dec/reset dyn=0, dec hold=0\n");
+
+ // CHECK-NEXT: structured{{.*}}
+ printf(" structured dec of dyn\n");
+ #pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+ {
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target data map(ompx_hold, tofrom: m) \
+ map(ompx_hold, alloc: r, d)
+ {
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target data map(tofrom: m) map(alloc: r, d)
+ {
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target data map(tofrom: m) map(alloc: r, d)
+ {
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ }
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ }
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ }
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ }
+ // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+ CHECK_PRESENCE(m, r, d);
+
+ // CHECK-NEXT: dynamic{{.*}}
+ printf(" dynamic dec/reset of dyn\n");
+ #pragma omp target enter data map(to: m) map(alloc: r, d)
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target enter data map(to: m) map(alloc: r, d)
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+ {
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target data map(ompx_hold, tofrom: m) \
+ map(ompx_hold, alloc: r, d)
+ {
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target exit data map(from: m) map(release: r)
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target exit data map(from: m) map(release: r) map(delete: d)
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ }
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target exit data map(from: m) map(release: r) map(delete: d)
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ }
+ // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+ CHECK_PRESENCE(m, r, d);
+
+ // -----------------------------------------------------------------------
+ // CHECK: check:{{.*}}
+ printf("check: dyn>0, hold>0, dec hold=0, dec/reset dyn=0\n");
+
+ // CHECK-NEXT: structured{{.*}}
+ printf(" structured dec of dyn\n");
+ #pragma omp target data map(tofrom: m) map(alloc: r, d)
+ {
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target data map(tofrom: m) map(alloc: r, d)
+ {
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target data map(ompx_hold, tofrom: m) \
+ map(ompx_hold, alloc: r, d)
+ {
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target data map(ompx_hold, tofrom: m) \
+ map(ompx_hold, alloc: r, d)
+ {
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ }
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ }
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ }
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ }
+ // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+ CHECK_PRESENCE(m, r, d);
+
+ // CHECK-NEXT: dynamic{{.*}}
+ printf(" dynamic dec/reset of dyn\n");
+ #pragma omp target enter data map(to: m) map(alloc: r, d)
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target enter data map(to: m) map(alloc: r, d)
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+ {
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target data map(ompx_hold, tofrom: m) \
+ map(ompx_hold, alloc: r, d)
+ {
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ }
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ }
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target exit data map(from: m) map(release: r)
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target exit data map(from: m) map(release: r) map(delete: d)
+ // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+ CHECK_PRESENCE(m, r, d);
+
+ return 0;
+}
diff --git a/openmp/libomptarget/test/mapping/ompx_hold/target.c b/openmp/libomptarget/test/mapping/ompx_hold/target.c
new file mode 100644
index 0000000000000..614970684aa3f
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/ompx_hold/target.c
@@ -0,0 +1,164 @@
+// RUN: %libomptarget-compile-generic -fopenmp-extensions
+// RUN: %libomptarget-run-generic | %fcheck-generic -strict-whitespace
+
+#include <omp.h>
+#include <stdio.h>
+
+#define CHECK_PRESENCE(Var1, Var2, Var3) \
+ printf(" presence of %s, %s, %s: %d, %d, %d\n", \
+ #Var1, #Var2, #Var3, \
+ omp_target_is_present(&Var1, omp_get_default_device()), \
+ omp_target_is_present(&Var2, omp_get_default_device()), \
+ omp_target_is_present(&Var3, omp_get_default_device()))
+
+int main() {
+ int m, r, d;
+ // CHECK: presence of m, r, d: 0, 0, 0
+ CHECK_PRESENCE(m, r, d);
+
+ // -----------------------------------------------------------------------
+ // CHECK-NEXT: check:{{.*}}
+ printf("check: dyn>0, hold=0, dec dyn=0\n");
+
+ // CHECK-NEXT: once
+ printf(" once\n");
+ #pragma omp target map(tofrom: m) map(alloc: r, d)
+ ;
+ // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+ CHECK_PRESENCE(m, r, d);
+
+ // CHECK-NEXT: twice
+ printf(" twice\n");
+ #pragma omp target data map(tofrom: m) map(alloc: r, d)
+ {
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target map(tofrom: m) map(alloc: r, d)
+ ;
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ }
+ // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+ CHECK_PRESENCE(m, r, d);
+
+ // -----------------------------------------------------------------------
+ // CHECK: check:{{.*}}
+ printf("check: dyn=0, hold>0, dec hold=0\n");
+
+ // CHECK-NEXT: once
+ printf(" once\n");
+ #pragma omp target map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+ ;
+ // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+ CHECK_PRESENCE(m, r, d);
+
+ // CHECK-NEXT: twice
+ printf(" twice\n");
+ #pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+ {
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+ ;
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ }
+ // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+ CHECK_PRESENCE(m, r, d);
+
+ // -----------------------------------------------------------------------
+ // CHECK: check:{{.*}}
+ printf("check: dyn>0, hold>0, dec dyn=0, dec hold=0\n");
+
+ // CHECK-NEXT: once each
+ printf(" once each\n");
+ #pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+ {
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target map(tofrom: m) map(alloc: r, d)
+ ;
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ }
+ // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+ CHECK_PRESENCE(m, r, d);
+
+ // CHECK-NEXT: twice each
+ printf(" twice each\n");
+ #pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+ {
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target data map(ompx_hold, tofrom: m) \
+ map(ompx_hold, alloc: r, d)
+ {
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target data map(tofrom: m) map(alloc: r, d)
+ {
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target map(tofrom: m) map(alloc: r, d)
+ ;
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ }
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ }
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ }
+ // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+ CHECK_PRESENCE(m, r, d);
+
+ // -----------------------------------------------------------------------
+ // CHECK: check:{{.*}}
+ printf("check: dyn>0, hold>0, dec hold=0, dec dyn=0\n");
+
+ // CHECK-NEXT: once each
+ printf(" once each\n");
+ #pragma omp target data map(tofrom: m) map(alloc: r, d)
+ {
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+ ;
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ }
+ // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+ CHECK_PRESENCE(m, r, d);
+
+ // CHECK-NEXT: twice each
+ printf(" twice each\n");
+ #pragma omp target data map(tofrom: m) map(alloc: r, d)
+ {
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target data map(tofrom: m) map(alloc: r, d)
+ {
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target data map(ompx_hold, tofrom: m) \
+ map(ompx_hold, alloc: r, d)
+ {
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ #pragma omp target map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
+ ;
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ }
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ }
+ // CHECK-NEXT: presence of m, r, d: 1, 1, 1
+ CHECK_PRESENCE(m, r, d);
+ }
+ // CHECK-NEXT: presence of m, r, d: 0, 0, 0
+ CHECK_PRESENCE(m, r, d);
+
+ return 0;
+}
diff --git a/openmp/libomptarget/test/offloading/info.c b/openmp/libomptarget/test/offloading/info.c
index b9635d547dd17..d935bbd2cc833 100644
--- a/openmp/libomptarget/test/offloading/info.c
+++ b/openmp/libomptarget/test/offloading/info.c
@@ -1,4 +1,7 @@
-// RUN: %libomptarget-compile-nvptx64-nvidia-cuda -gline-tables-only && env LIBOMPTARGET_INFO=63 %libomptarget-run-nvptx64-nvidia-cuda 2>&1 | %fcheck-nvptx64-nvidia-cuda -allow-empty -check-prefix=INFO
+// RUN: %libomptarget-compile-nvptx64-nvidia-cuda \
+// RUN: -gline-tables-only -fopenmp-extensions
+// RUN: env LIBOMPTARGET_INFO=63 %libomptarget-run-nvptx64-nvidia-cuda 2>&1 | \
+// RUN: %fcheck-nvptx64-nvidia-cuda -allow-empty -check-prefix=INFO
// REQUIRES: nvptx64-nvidia-cuda
#include <stdio.h>
@@ -23,24 +26,24 @@ int main() {
// INFO: Libomptarget device 0 info: alloc(A[0:64])[256]
// INFO: Libomptarget device 0 info: tofrom(B[0:64])[256]
// INFO: Libomptarget device 0 info: to(C[0:64])[256]
-// INFO: Libomptarget device 0 info: Creating new map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, RefCount=1, Name=A[0:64]
-// INFO: Libomptarget device 0 info: Creating new map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, RefCount=1, Name=B[0:64]
+// INFO: Libomptarget device 0 info: Creating new map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, DynRefCount=1, HoldRefCount=0, Name=A[0:64]
+// INFO: Libomptarget device 0 info: Creating new map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, DynRefCount=0, HoldRefCount=1, Name=B[0:64]
// INFO: Libomptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=256, Name=B[0:64]
-// INFO: Libomptarget device 0 info: Creating new map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, RefCount=1, Name=C[0:64]
+// INFO: Libomptarget device 0 info: Creating new map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, DynRefCount=1, HoldRefCount=0, Name=C[0:64]
// INFO: Libomptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=256, Name=C[0:64]
// INFO: Libomptarget device 0 info: OpenMP Host-Device pointer mappings after block at info.c:{{[0-9]+}}:{{[0-9]+}}:
-// INFO: Libomptarget device 0 info: Host Ptr Target Ptr Size (B) RefCount Declaration
-// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 1 C[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
-// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 1 B[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
-// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 1 A[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
+// INFO: Libomptarget device 0 info: Host Ptr Target Ptr Size (B) DynRefCount HoldRefCount Declaration
+// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 1 0 C[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
+// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 0 1 B[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
+// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 1 0 A[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
// INFO: Libomptarget device 0 info: Entering OpenMP kernel at info.c:{{[0-9]+}}:{{[0-9]+}} with 1 arguments:
// INFO: Libomptarget device 0 info: firstprivate(val)[4]
// INFO: CUDA device 0 info: Launching kernel __omp_offloading_{{.*}}main{{.*}} with {{[0-9]+}} blocks and {{[0-9]+}} threads in Generic mode
// INFO: Libomptarget device 0 info: OpenMP Host-Device pointer mappings after block at info.c:{{[0-9]+}}:{{[0-9]+}}:
-// INFO: Libomptarget device 0 info: Host Ptr Target Ptr Size (B) RefCount Declaration
-// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 1 C[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
-// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 1 B[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
-// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 1 A[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
+// INFO: Libomptarget device 0 info: Host Ptr Target Ptr Size (B) DynRefCount HoldRefCount Declaration
+// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 1 0 C[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
+// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 0 1 B[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
+// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 1 0 A[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
// INFO: Libomptarget device 0 info: Exiting OpenMP data region at info.c:{{[0-9]+}}:{{[0-9]+}} with 3 arguments:
// INFO: Libomptarget device 0 info: alloc(A[0:64])[256]
// INFO: Libomptarget device 0 info: tofrom(B[0:64])[256]
@@ -50,9 +53,9 @@ int main() {
// INFO: Libomptarget device 0 info: Removing map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, Name=B[0:64]
// INFO: Libomptarget device 0 info: Removing map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, Name=A[0:64]
// INFO: Libomptarget device 0 info: OpenMP Host-Device pointer mappings after block at info.c:[[#%u,]]:[[#%u,]]:
-// INFO: Libomptarget device 0 info: Host Ptr Target Ptr Size (B) RefCount Declaration
-// INFO: Libomptarget device 0 info: [[#%#x,]] [[#%#x,]] 4 INF unknown at unknown:0:0
-#pragma omp target data map(alloc:A[0:N]) map(tofrom:B[0:N]) map(to:C[0:N])
+// INFO: Libomptarget device 0 info: Host Ptr Target Ptr Size (B) DynRefCount HoldRefCount Declaration
+// INFO: Libomptarget device 0 info: [[#%#x,]] [[#%#x,]] 4 INF 0 unknown at unknown:0:0
+#pragma omp target data map(alloc:A[0:N]) map(ompx_hold,tofrom:B[0:N]) map(to:C[0:N])
#pragma omp target firstprivate(val)
{ val = 1; }
More information about the Openmp-commits
mailing list