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 @@ -173,7 +173,7 @@ auto &HT = *lr.Entry; IsNew = false; - if (UpdateRefCount) + if (!CONSIDERED_INF(HT.RefCount) && UpdateRefCount) ++HT.RefCount; uintptr_t tp = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin); @@ -231,7 +231,7 @@ auto &HT = *lr.Entry; IsLast = !(HT.RefCount > 1); - if (HT.RefCount > 1 && UpdateRefCount) + if (!CONSIDERED_INF(HT.RefCount) && HT.RefCount > 1 && UpdateRefCount) --HT.RefCount; uintptr_t tp = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin); @@ -279,9 +279,9 @@ LookupResult lr = lookupMapping(HstPtrBegin, Size); if (lr.Flags.IsContained || lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) { auto &HT = *lr.Entry; - if (ForceDelete) + if (ForceDelete && !CONSIDERED_INF(HT.RefCount)) HT.RefCount = 1; - if (--HT.RefCount <= 0) { + if (!CONSIDERED_INF(HT.RefCount) && --HT.RefCount <= 0) { assert(HT.RefCount == 0 && "did not expect a negative ref count"); DP("Deleting tgt data " DPxMOD " of size %ld\n", DPxPTR(HT.TgtPtrBegin), Size); 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; +}