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,9 @@ 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}; /// 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 @@ -124,8 +124,10 @@ TargetPointerResultTy TPR = Device.getTgtPtrBegin(const_cast(Ptr), 1, IsLast, /*UpdateRefCount=*/false, - /*UseHoldRefCount=*/false, IsHostPtr); - int Rc = (TPR.TargetPointer != NULL); + /*UseHoldRefCount=*/false, IsHostPtr, + /*MustContain=*/false, + /*ForceDelete=*/false); + int Rc = TPR.Flags.IsPresent; // 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 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 @@ -365,6 +365,7 @@ void *TargetPointer = NULL; bool IsNew = false; + bool IsPresent = true; IsHostPtr = false; IsLast = false; LookupResult LR = lookupMapping(HDTTMap, HstPtrBegin, Size); @@ -416,11 +417,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.Flags.IsPresent && !TPR.Flags.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.Flags.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.Flags.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.Flags.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 = 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()); +}