diff --git a/openmp/libomptarget/include/device.h b/openmp/libomptarget/include/device.h --- a/openmp/libomptarget/include/device.h +++ b/openmp/libomptarget/include/device.h @@ -281,7 +281,13 @@ unsigned IsNewEntry : 1; /// If the pointer is actually a host pointer (when unified memory enabled) unsigned IsHostPointer : 1; - } Flags = {0, 0}; + /// If the pointer is present in the mapping table. + unsigned IsPresent : 1; + } Flags = {0, 0, 0}; + + bool isPresent() const { return Flags.IsPresent; } + + bool isHostPointer() const { return Flags.IsHostPointer; } /// The corresponding map table entry which is stable. HostDataToTargetTy *Entry = nullptr; diff --git a/openmp/libomptarget/src/api.cpp b/openmp/libomptarget/src/api.cpp --- a/openmp/libomptarget/src/api.cpp +++ b/openmp/libomptarget/src/api.cpp @@ -125,13 +125,7 @@ Device.getTgtPtrBegin(const_cast(Ptr), 1, IsLast, /*UpdateRefCount=*/false, /*UseHoldRefCount=*/false, IsHostPtr); - int Rc = (TPR.TargetPointer != NULL); - // Under unified memory the host pointer can be returned by the - // getTgtPtrBegin() function which means that there is no device - // corresponding point for ptr. This function should return false - // in that situation. - if (PM->RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) - Rc = !IsHostPtr; + int Rc = TPR.isPresent(); DP("Call to omp_target_is_present returns %d\n", Rc); return Rc; } 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 @@ -216,6 +216,7 @@ void *TargetPointer = nullptr; bool IsHostPtr = false; + bool IsPresent = true; bool IsNew = false; LookupResult LR = lookupMapping(HDTTMap, HstPtrBegin, Size); @@ -275,6 +276,7 @@ DP("Return HstPtrBegin " DPxMOD " Size=%" PRId64 " for unified shared " "memory\n", DPxPTR((uintptr_t)HstPtrBegin), Size); + IsPresent = false; IsHostPtr = true; TargetPointer = HstPtrBegin; } @@ -303,6 +305,9 @@ Entry->dynRefCountToStr().c_str(), Entry->holdRefCountToStr().c_str(), (HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "unknown"); TargetPointer = (void *)Ptr; + } else { + // This entry is not present and we did not create a new entry for it. + IsPresent = false; } // If the target pointer is valid, and we need to transfer data, issue the @@ -351,7 +356,7 @@ } } - return {{IsNew, IsHostPtr}, Entry, TargetPointer}; + return {{IsNew, IsHostPtr, IsPresent}, Entry, TargetPointer}; } // Used by targetDataBegin, targetDataEnd, targetDataUpdate and target. @@ -365,6 +370,7 @@ void *TargetPointer = NULL; bool IsNew = false; + bool IsPresent = true; IsHostPtr = false; IsLast = false; LookupResult LR = lookupMapping(HDTTMap, HstPtrBegin, Size); @@ -416,11 +422,18 @@ DP("Get HstPtrBegin " DPxMOD " Size=%" PRId64 " for unified shared " "memory\n", DPxPTR((uintptr_t)HstPtrBegin), Size); + IsPresent = false; IsHostPtr = true; TargetPointer = HstPtrBegin; + } else { + // OpenMP Specification v5.2: if a matching list item is not found, the + // pointer retains its original value as per firstprivate semantics. + IsPresent = false; + IsHostPtr = false; + TargetPointer = HstPtrBegin; } - return {{IsNew, IsHostPtr}, LR.Entry, TargetPointer}; + return {{IsNew, IsHostPtr, IsPresent}, LR.Entry, TargetPointer}; } // Return the target pointer begin (where the data will be moved). 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 @@ -744,7 +744,8 @@ HstPtrBegin, DataSize, IsLast, UpdateRef, HasHoldModifier, IsHostPtr, !IsImplicit, ForceDelete); void *TgtPtrBegin = TPR.TargetPointer; - if (!TgtPtrBegin && (DataSize || HasPresentModifier)) { + if (!TPR.isPresent() && !TPR.isHostPointer() && + (DataSize || HasPresentModifier)) { DP("Mapping does not exist (%s)\n", (HasPresentModifier ? "'present' map type modifier" : "ignored")); if (HasPresentModifier) { @@ -779,7 +780,7 @@ // construct and a corresponding list item of the original list item is not // present in the device data environment on exit from the region then the // list item is ignored." - if (!TgtPtrBegin) + if (!TPR.isPresent()) continue; bool DelEntry = IsLast; @@ -921,7 +922,7 @@ HstPtrBegin, ArgSize, IsLast, /*UpdateRefCount=*/false, /*UseHoldRefCount=*/false, IsHostPtr, /*MustContain=*/true); void *TgtPtrBegin = TPR.TargetPointer; - if (!TgtPtrBegin) { + if (!TPR.isPresent()) { DP("hst data:" DPxMOD " not found, becomes a noop\n", DPxPTR(HstPtrBegin)); if (ArgType & OMP_TGT_MAPTYPE_PRESENT) { MESSAGE("device mapping required by 'present' motion modifier does not " @@ -1349,7 +1350,7 @@ HstPtrVal, ArgSizes[I], IsLast, /*UpdateRefCount=*/false, /*UseHoldRefCount=*/false, IsHostPtr); PointerTgtPtrBegin = TPR.TargetPointer; - if (!PointerTgtPtrBegin) { + if (!TPR.isPresent()) { DP("No lambda captured variable mapped (" DPxMOD ") - ignored\n", DPxPTR(HstPtrVal)); continue; diff --git a/openmp/libomptarget/test/mapping/implicit_device_ptr.c b/openmp/libomptarget/test/mapping/implicit_device_ptr.c new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/mapping/implicit_device_ptr.c @@ -0,0 +1,26 @@ +// RUN: %libomptarget-compile-run-and-check-generic + +#include +#include + +// OpenMP 5.1. sec 5.8.6 "Pointer Initialization for Device Data Environments" +// p. 160 L32-33: "If a matching mapped list item is not found, the pointer +// retains its original value as per the32 firstprivate semantics described in +// Section 5.4.4." + +int main(void) { + int *A = (int *)omp_target_alloc(sizeof(int), omp_get_default_device()); + +#pragma omp target + { *A = 1; } + + int Result = 0; +#pragma omp target map(from : Result) + { Result = *A; } + + // CHECK: PASS + if (Result == 1) + printf("PASS\n"); + + omp_target_free(A, omp_get_default_device()); +}