Index: openmp/libomptarget/include/omptarget.h =================================================================== --- openmp/libomptarget/include/omptarget.h +++ openmp/libomptarget/include/omptarget.h @@ -56,6 +56,10 @@ 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 Index: openmp/libomptarget/src/api.cpp =================================================================== --- openmp/libomptarget/src/api.cpp +++ openmp/libomptarget/src/api.cpp @@ -106,7 +106,8 @@ bool IsLast; // not used bool IsHostPtr; void *TgtPtr = Device.getTgtPtrBegin(const_cast(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 Index: openmp/libomptarget/src/device.h =================================================================== --- openmp/libomptarget/src/device.h +++ openmp/libomptarget/src/device.h @@ -50,55 +50,108 @@ 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); + } 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) {} + + /// 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; + } - uint64_t getRefCount() const { return RefCount; } + /// 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; } }; @@ -195,18 +248,18 @@ bool isDataExchangable(const DeviceTy &DstDevice); LookupResult lookupMapping(void *HstPtrBegin, int64_t Size); - TargetPointerResultTy getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase, - int64_t Size, - map_var_info_t HstPtrName, - bool IsImplicit, bool UpdateRefCount, - bool HasCloseModifier, - bool HasPresentModifier); + TargetPointerResultTy + getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase, int64_t Size, + map_var_info_t HstPtrName, bool IsImplicit, + bool UpdateRefCount, bool HasCloseModifier, + bool HasPresentModifier, bool HasHoldModifier); 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); Index: openmp/libomptarget/src/device.cpp =================================================================== --- openmp/libomptarget/src/device.cpp +++ openmp/libomptarget/src/device.cpp @@ -82,14 +82,17 @@ /*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()); + HostDataToTargetMap.insert(newEntry); DataMapMtx.unlock(); @@ -102,7 +105,13 @@ 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(); @@ -111,11 +120,12 @@ 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; } @@ -180,7 +190,7 @@ DeviceTy::getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase, int64_t Size, map_var_info_t HstPtrName, bool IsImplicit, bool UpdateRefCount, bool HasCloseModifier, - bool HasPresentModifier) { + bool HasPresentModifier, bool HasHoldModifier) { void *TargetPointer = NULL; bool IsNew = false; bool IsHostPtr = false; @@ -195,21 +205,27 @@ 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) + IsNew = false; + 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) { @@ -252,13 +268,15 @@ 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; } @@ -271,8 +289,9 @@ // 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; @@ -282,35 +301,39 @@ 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 @@ -342,7 +365,7 @@ } 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; @@ -352,7 +375,7 @@ 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); Index: openmp/libomptarget/src/omptarget.cpp =================================================================== --- openmp/libomptarget/src/omptarget.cpp +++ openmp/libomptarget/src/omptarget.cpp @@ -157,7 +157,8 @@ (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 @@ // 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 @@ -489,7 +491,7 @@ // when HasPresentModifier. Pointer_TPR = Device.getOrAllocTgtPtr( HstPtrBase, HstPtrBase, sizeof(void *), nullptr, IsImplicit, - UpdateRef, HasCloseModifier, HasPresentModifier); + UpdateRef, HasCloseModifier, HasPresentModifier, HasHoldModifier); PointerTgtPtrBegin = Pointer_TPR.TargetPointer; IsHostPtr = Pointer_TPR.Flags.IsHostPointer; if (!PointerTgtPtrBegin) { @@ -511,9 +513,9 @@ (!FromMapper || i != 0); // subsequently update ref count of pointee } - auto TPR = Device.getOrAllocTgtPtr(HstPtrBegin, HstPtrBase, data_size, - HstPtrName, IsImplicit, UpdateRef, - HasCloseModifier, HasPresentModifier); + auto TPR = Device.getOrAllocTgtPtr( + HstPtrBegin, HstPtrBase, data_size, HstPtrName, IsImplicit, UpdateRef, + HasCloseModifier, HasPresentModifier, HasHoldModifier); void *TgtPtrBegin = TPR.TargetPointer; IsHostPtr = TPR.Flags.IsHostPointer; // If data_size==0, then the argument could be a zero-length pointer to @@ -612,10 +614,13 @@ 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 @@ -682,11 +687,12 @@ 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")); @@ -804,7 +810,8 @@ // Add pointer to the buffer for later deallocation if (DelEntry) - DeallocTgtPtrs.emplace_back(HstPtrBegin, DataSize, HasCloseModifier); + DeallocTgtPtrs.emplace_back(HstPtrBegin, DataSize, HasCloseModifier, + HasHoldModifier); } } @@ -821,7 +828,7 @@ 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; @@ -836,8 +843,9 @@ 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) { @@ -1296,8 +1304,9 @@ 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)); @@ -1353,7 +1362,8 @@ if (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ) HstPtrBase = *reinterpret_cast(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); Index: openmp/libomptarget/src/private.h =================================================================== --- openmp/libomptarget/src/private.h +++ openmp/libomptarget/src/private.h @@ -112,16 +112,18 @@ 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(); Index: openmp/libomptarget/test/mapping/ompx_hold/omp_target_disassociate_ptr.c =================================================================== --- /dev/null +++ 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 +#include +#include +#include + +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; +} Index: openmp/libomptarget/test/mapping/ompx_hold/struct.c =================================================================== --- /dev/null +++ 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 +#include + +#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; +} Index: openmp/libomptarget/test/mapping/ompx_hold/target-data.c =================================================================== --- /dev/null +++ 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 +#include + +#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; +} Index: openmp/libomptarget/test/mapping/ompx_hold/target.c =================================================================== --- /dev/null +++ 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 +#include + +#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; +} Index: openmp/libomptarget/test/offloading/info.c =================================================================== --- openmp/libomptarget/test/offloading/info.c +++ 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 @@ -23,24 +26,24 @@ // 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 @@ // 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; }