[Openmp-commits] [openmp] e244145 - [LIBOMPTARGET] Do not increment/decrement the refcount for "declare target" objects
George Rokos via Openmp-commits
openmp-commits at lists.llvm.org
Tue Jan 14 16:31:30 PST 2020
Author: George Rokos
Date: 2020-01-14T16:30:38-08:00
New Revision: e244145ab08ae79ea3d22c2fe479ec084dbd7742
URL: https://github.com/llvm/llvm-project/commit/e244145ab08ae79ea3d22c2fe479ec084dbd7742
DIFF: https://github.com/llvm/llvm-project/commit/e244145ab08ae79ea3d22c2fe479ec084dbd7742.diff
LOG: [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
Added:
openmp/libomptarget/test/mapping/delete_inf_refcount.c
Modified:
openmp/libomptarget/src/device.cpp
openmp/libomptarget/src/device.h
openmp/libomptarget/src/omptarget.cpp
Removed:
################################################################################
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 <cassert>
#include <cstddef>
-#include <climits>
#include <list>
#include <map>
#include <mutex>
@@ -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<HostDataToTargetTy> 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 <stdio.h>
+#include <omp.h>
+
+#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;
+}
More information about the Openmp-commits
mailing list