From e244145ab08ae79ea3d22c2fe479ec084dbd7742 Mon Sep 17 00:00:00 2001 From: George Rokos Date: Tue, 14 Jan 2020 16:30:38 -0800 Subject: [PATCH] [LIBOMPTARGET] Do not increment/decrement the refcount for "declare target" objects The reference counter for global objects marked with declare target is INF. This patch prevents the runtime from incrementing /decrementing INF refcounts. Without it, the map(delete: global_object) directive actually deallocates the global on the device. With this patch, such a directive becomes a no-op. Differential Revision: https://reviews.llvm.org/D72525 --- openmp/libomptarget/src/device.cpp | 45 +++++++-------- openmp/libomptarget/src/device.h | 55 ++++++++++++++----- openmp/libomptarget/src/omptarget.cpp | 6 +- .../test/mapping/delete_inf_refcount.c | 32 +++++++++++ 4 files changed, 95 insertions(+), 43 deletions(-) create mode 100644 openmp/libomptarget/test/mapping/delete_inf_refcount.c diff --git a/openmp/libomptarget/src/device.cpp b/openmp/libomptarget/src/device.cpp index cf7e0fe0c1b2..41a1b53de1f9 100644 --- a/openmp/libomptarget/src/device.cpp +++ b/openmp/libomptarget/src/device.cpp @@ -44,16 +44,12 @@ int DeviceTy::associatePtr(void *HstPtrBegin, void *TgtPtrBegin, int64_t Size) { } } - // 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 @@ int DeviceTy::disassociatePtr(void *HstPtrBegin) { 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 @@ int DeviceTy::disassociatePtr(void *HstPtrBegin) { } // 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 @@ void *DeviceTy::getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase, 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 @@ void *DeviceTy::getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast, 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 @@ int DeviceTy::deallocTgtPtr(void *HstPtrBegin, int64_t Size, bool ForceDelete, 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/device.h b/openmp/libomptarget/src/device.h index d33512bb08e7..8379f0c65ae4 100644 --- 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,9 +25,6 @@ struct RTLInfoTy; 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 { uintptr_t HstPtrBase; // host info. @@ -36,18 +33,48 @@ struct HostDataToTargetTy { uintptr_t TgtPtrBegin; // target info. - long RefCount; +private: + uint64_t RefCount; + 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) {} + + uint64_t getRefCount() const { + return RefCount; + } + + uint64_t resetRefCount() { + if (RefCount != INFRefCount) + RefCount = 1; + + return RefCount; + } + + uint64_t incRefCount() { + if (RefCount != INFRefCount) { + ++RefCount; + assert(RefCount < INFRefCount && "refcount overflow"); + } + + return RefCount; + } + + uint64_t decRefCount() { + if (RefCount != INFRefCount) { + assert(RefCount > 0 && "refcount underflow"); + --RefCount; + } + + return RefCount; + } + + bool isRefCountInf() const { + return RefCount == INFRefCount; + } }; typedef std::list HostDataToTargetListTy; @@ -129,7 +156,7 @@ struct DeviceTy { 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/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp index 2feb7c89f41e..fed7dcc189f0 100644 --- a/openmp/libomptarget/src/omptarget.cpp +++ b/openmp/libomptarget/src/omptarget.cpp @@ -139,7 +139,7 @@ static int InitLibrary(DeviceTy& Device) { (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 @@ int target_data_begin(DeviceTy &Device, int32_t arg_num, } 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 @@ int target_data_end(DeviceTy &Device, int32_t arg_num, void **args_base, !(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 index 000000000000..b4106be04ab7 --- /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; +}