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 @@ -25,29 +25,51 @@ 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. +private: long RefCount; + bool INFRefCount; - 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 INF = false) : HstPtrBase(BP), HstPtrBegin(B), HstPtrEnd(E), - TgtPtrBegin(TB), RefCount(RF) {} + TgtPtrBegin(TB), RefCount(INF ? LONG_MAX : 1), INFRefCount(INF) {} + + long getRefCount() const { + return RefCount; + } + + void setRefCount(long Value) { + if (!INFRefCount) + RefCount = Value; + } + + long incRefCount() { + if (!INFRefCount) + ++RefCount; + + return RefCount; + } + + long decRefCount() { + if (!INFRefCount) + --RefCount; + + return RefCount; + } + + bool isRefCountInf() const { + return INFRefCount; + } }; typedef std::list HostDataToTargetListTy; 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 /*INFRefCount*/); 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(); @@ -102,7 +98,7 @@ 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; } } @@ -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,9 @@ 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.setRefCount(1); + if (HT.decRefCount() <= 0) { + assert(HT.getRefCount() == 0 && "did not expect a negative ref count"); 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 /*INFRefCount*/)); } } Device.DataMapMtx.unlock(); 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; +}