diff --git a/openmp/libomptarget/src/device.h b/openmp/libomptarget/src/device.h --- a/openmp/libomptarget/src/device.h +++ b/openmp/libomptarget/src/device.h @@ -13,8 +13,8 @@ #ifndef _OMPTARGET_DEVICE_H #define _OMPTARGET_DEVICE_H +#include #include -#include #include #include #include @@ -25,29 +25,58 @@ struct __tgt_bin_desc; struct __tgt_target_table; -#define INF_REF_CNT (LONG_MAX>>1) // leave room for additions/subtractions -#define CONSIDERED_INF(x) (x > (INF_REF_CNT>>1)) - /// Map between host data and target data. struct HostDataToTargetTy { +public: uintptr_t HstPtrBase; // host info. uintptr_t HstPtrBegin; uintptr_t HstPtrEnd; // non-inclusive. uintptr_t TgtPtrBegin; // target info. - long RefCount; +private: + uint64_t RefCount; + bool IsRefCountINF; + static const uint64_t INFRefCount = ~(uint64_t)0; - HostDataToTargetTy() - : HstPtrBase(0), HstPtrBegin(0), HstPtrEnd(0), - TgtPtrBegin(0), RefCount(0) {} - HostDataToTargetTy(uintptr_t BP, uintptr_t B, uintptr_t E, uintptr_t TB) - : HstPtrBase(BP), HstPtrBegin(B), HstPtrEnd(E), - TgtPtrBegin(TB), RefCount(1) {} +public: HostDataToTargetTy(uintptr_t BP, uintptr_t B, uintptr_t E, uintptr_t TB, - long RF) + bool IsINF = false) : HstPtrBase(BP), HstPtrBegin(B), HstPtrEnd(E), - TgtPtrBegin(TB), RefCount(RF) {} + TgtPtrBegin(TB), RefCount(IsINF ? INFRefCount : 1), IsRefCountINF(IsINF) {} + + uint64_t getRefCount() const { + return RefCount; + } + + uint64_t resetRefCount() { + if (!IsRefCountINF) + RefCount = 1; + + return RefCount; + } + + uint64_t incRefCount() { + if (!IsRefCountINF) { + assert(RefCount < INFRefCount && "refcount overflow"); + ++RefCount; + } + + return RefCount; + } + + uint64_t decRefCount() { + if (!IsRefCountINF) { + assert(RefCount > 0 && "refcount underflow"); + --RefCount; + } + + return RefCount; + } + + bool isRefCountInf() const { + return IsRefCountINF; + } }; typedef std::list HostDataToTargetListTy; @@ -129,7 +158,7 @@ return *this; } - long getMapEntryRefCnt(void *HstPtrBegin); + uint64_t getMapEntryRefCnt(void *HstPtrBegin); LookupResult lookupMapping(void *HstPtrBegin, int64_t Size); void *getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase, int64_t Size, bool &IsNew, bool &IsHostPtr, bool IsImplicit, bool UpdateRefCount = true, diff --git a/openmp/libomptarget/src/device.cpp b/openmp/libomptarget/src/device.cpp --- a/openmp/libomptarget/src/device.cpp +++ b/openmp/libomptarget/src/device.cpp @@ -44,16 +44,12 @@ } } - // Mapping does not exist, allocate it - HostDataToTargetTy newEntry; - - // Set up missing fields - newEntry.HstPtrBase = (uintptr_t) HstPtrBegin; - newEntry.HstPtrBegin = (uintptr_t) HstPtrBegin; - newEntry.HstPtrEnd = (uintptr_t) HstPtrBegin + Size; - newEntry.TgtPtrBegin = (uintptr_t) TgtPtrBegin; - // refCount must be infinite - newEntry.RefCount = INF_REF_CNT; + // Mapping does not exist, allocate it with refCount=INF + HostDataToTargetTy newEntry((uintptr_t) HstPtrBegin /*HstPtrBase*/, + (uintptr_t) HstPtrBegin /*HstPtrBegin*/, + (uintptr_t) HstPtrBegin + Size /*HstPtrEnd*/, + (uintptr_t) TgtPtrBegin /*TgtPtrBegin*/, + true /*IsRefCountINF*/); DP("Creating new map entry: HstBase=" DPxMOD ", HstBegin=" DPxMOD ", HstEnd=" DPxMOD ", TgtBegin=" DPxMOD "\n", DPxPTR(newEntry.HstPtrBase), @@ -74,7 +70,7 @@ ii != HostDataToTargetMap.end(); ++ii) { if ((uintptr_t)HstPtrBegin == ii->HstPtrBegin) { // Mapping exists - if (CONSIDERED_INF(ii->RefCount)) { + if (ii->isRefCountInf()) { DP("Association found, removing it\n"); HostDataToTargetMap.erase(ii); DataMapMtx.unlock(); @@ -94,21 +90,21 @@ } // Get ref count of map entry containing HstPtrBegin -long DeviceTy::getMapEntryRefCnt(void *HstPtrBegin) { +uint64_t DeviceTy::getMapEntryRefCnt(void *HstPtrBegin) { uintptr_t hp = (uintptr_t)HstPtrBegin; - long RefCnt = -1; + uint64_t RefCnt = 0; DataMapMtx.lock(); for (auto &HT : HostDataToTargetMap) { if (hp >= HT.HstPtrBegin && hp < HT.HstPtrEnd) { DP("DeviceTy::getMapEntry: requested entry found\n"); - RefCnt = HT.RefCount; + RefCnt = HT.getRefCount(); break; } } DataMapMtx.unlock(); - if (RefCnt < 0) { + if (RefCnt == 0) { DP("DeviceTy::getMapEntry: requested entry not found\n"); } @@ -174,15 +170,14 @@ IsNew = false; if (UpdateRefCount) - ++HT.RefCount; + HT.incRefCount(); uintptr_t tp = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin); DP("Mapping exists%s with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", " "Size=%ld,%s RefCount=%s\n", (IsImplicit ? " (implicit)" : ""), DPxPTR(HstPtrBegin), DPxPTR(tp), Size, (UpdateRefCount ? " updated" : ""), - (CONSIDERED_INF(HT.RefCount)) ? "INF" : - std::to_string(HT.RefCount).c_str()); + HT.isRefCountInf() ? "INF" : std::to_string(HT.getRefCount()).c_str()); rc = (void *)tp; } else if ((lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) && !IsImplicit) { // Explicit extension of mapped data - not allowed. @@ -229,17 +224,16 @@ if (lr.Flags.IsContained || lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) { auto &HT = *lr.Entry; - IsLast = !(HT.RefCount > 1); + IsLast = HT.getRefCount() == 1; - if (HT.RefCount > 1 && UpdateRefCount) - --HT.RefCount; + if (!IsLast && UpdateRefCount) + HT.decRefCount(); uintptr_t tp = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin); DP("Mapping exists with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", " "Size=%ld,%s RefCount=%s\n", DPxPTR(HstPtrBegin), DPxPTR(tp), Size, (UpdateRefCount ? " updated" : ""), - (CONSIDERED_INF(HT.RefCount)) ? "INF" : - std::to_string(HT.RefCount).c_str()); + HT.isRefCountInf() ? "INF" : std::to_string(HT.getRefCount()).c_str()); rc = (void *)tp; } else if (RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) { // If the value isn't found in the mapping and unified shared memory @@ -280,9 +274,8 @@ if (lr.Flags.IsContained || lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) { auto &HT = *lr.Entry; if (ForceDelete) - HT.RefCount = 1; - if (--HT.RefCount <= 0) { - assert(HT.RefCount == 0 && "did not expect a negative ref count"); + HT.resetRefCount(); + if (HT.decRefCount() == 0) { DP("Deleting tgt data " DPxMOD " of size %ld\n", DPxPTR(HT.TgtPtrBegin), Size); RTL->data_delete(RTLDeviceID, (void *)HT.TgtPtrBegin); diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp --- a/openmp/libomptarget/src/omptarget.cpp +++ b/openmp/libomptarget/src/omptarget.cpp @@ -139,7 +139,7 @@ (uintptr_t)CurrHostEntry->addr /*HstPtrBegin*/, (uintptr_t)CurrHostEntry->addr + CurrHostEntry->size /*HstPtrEnd*/, (uintptr_t)CurrDeviceEntry->addr /*TgtPtrBegin*/, - INF_REF_CNT /*RefCount*/)); + true /*IsRefCountINF*/)); } } Device.DataMapMtx.unlock(); @@ -301,7 +301,7 @@ } else if (arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF) { // Copy data only if the "parent" struct has RefCount==1. int32_t parent_idx = member_of(arg_types[i]); - long parent_rc = Device.getMapEntryRefCnt(args[parent_idx]); + uint64_t parent_rc = Device.getMapEntryRefCnt(args[parent_idx]); assert(parent_rc > 0 && "parent struct not found"); if (parent_rc == 1) { copy = true; @@ -402,7 +402,7 @@ !(arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)) { // Copy data only if the "parent" struct has RefCount==1. int32_t parent_idx = member_of(arg_types[i]); - long parent_rc = Device.getMapEntryRefCnt(args[parent_idx]); + uint64_t parent_rc = Device.getMapEntryRefCnt(args[parent_idx]); assert(parent_rc > 0 && "parent struct not found"); if (parent_rc == 1) { CopyMember = true; diff --git a/openmp/libomptarget/test/mapping/delete_inf_refcount.c b/openmp/libomptarget/test/mapping/delete_inf_refcount.c new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/mapping/delete_inf_refcount.c @@ -0,0 +1,32 @@ +// RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu +// RUN: %libomptarget-compile-run-and-check-powerpc64-ibm-linux-gnu +// RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu +// RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu + +#include +#include + +#pragma omp declare target +int isHost; +#pragma omp end declare target + +int main(void) { + isHost = -1; + +#pragma omp target enter data map(to: isHost) + +#pragma omp target + { isHost = omp_is_initial_device(); } +#pragma omp target update from(isHost) + + if (isHost < 0) { + printf("Runtime error, isHost=%d\n", isHost); + } + +#pragma omp target exit data map(delete: isHost) + + // CHECK: Target region executed on the device + printf("Target region executed on the %s\n", isHost ? "host" : "device"); + + return isHost; +}