[Openmp-commits] [openmp] 6e8d93e - [Libomptarget] Implement OpenMP 5.2 semantics for device pointers
Joseph Huber via Openmp-commits
openmp-commits at lists.llvm.org
Wed Sep 7 15:01:27 PDT 2022
Author: Joseph Huber
Date: 2022-09-07T17:01:14-05:00
New Revision: 6e8d93e5c2351efcf3f6fd658af642b027da3bc4
URL: https://github.com/llvm/llvm-project/commit/6e8d93e5c2351efcf3f6fd658af642b027da3bc4
DIFF: https://github.com/llvm/llvm-project/commit/6e8d93e5c2351efcf3f6fd658af642b027da3bc4.diff
LOG: [Libomptarget] Implement OpenMP 5.2 semantics for device pointers
In OpenMP 5.2, ยง5.8.6, page 160 line 32-33, when a device pointer
allocated by omp_target_alloc has implicitly been included on a target
construct as a zero-length array, the pointer initialisation should not
find a matching mapped list item, and so should retain its value as a
firstprivate variable. Previously, we would return a null pointer if the
list item was not found. This patch updates the map handling to the
OpenMP 5.2 semantics.
Reviewed By: jdoerfert, ye-luo
Differential Revision: https://reviews.llvm.org/D133447
Added:
openmp/libomptarget/test/mapping/implicit_device_ptr.c
Modified:
openmp/libomptarget/include/device.h
openmp/libomptarget/src/api.cpp
openmp/libomptarget/src/device.cpp
openmp/libomptarget/src/omptarget.cpp
Removed:
################################################################################
diff --git a/openmp/libomptarget/include/device.h b/openmp/libomptarget/include/device.h
index 2e4ebf5846f8e..5141899a19fa9 100644
--- a/openmp/libomptarget/include/device.h
+++ b/openmp/libomptarget/include/device.h
@@ -281,7 +281,13 @@ struct TargetPointerResultTy {
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
index ef54511ef8dd2..10bf242db0f99 100644
--- a/openmp/libomptarget/src/api.cpp
+++ b/openmp/libomptarget/src/api.cpp
@@ -125,13 +125,7 @@ EXTERN int omp_target_is_present(const void *Ptr, int DeviceNum) {
Device.getTgtPtrBegin(const_cast<void *>(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
index 29b5a6c7b9e62..43e9d13a6e3ce 100644
--- a/openmp/libomptarget/src/device.cpp
+++ b/openmp/libomptarget/src/device.cpp
@@ -216,6 +216,7 @@ TargetPointerResultTy DeviceTy::getTargetPointer(
void *TargetPointer = nullptr;
bool IsHostPtr = false;
+ bool IsPresent = true;
bool IsNew = false;
LookupResult LR = lookupMapping(HDTTMap, HstPtrBegin, Size);
@@ -275,6 +276,7 @@ TargetPointerResultTy DeviceTy::getTargetPointer(
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 @@ TargetPointerResultTy DeviceTy::getTargetPointer(
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 @@ TargetPointerResultTy DeviceTy::getTargetPointer(
}
}
- return {{IsNew, IsHostPtr}, Entry, TargetPointer};
+ return {{IsNew, IsHostPtr, IsPresent}, Entry, TargetPointer};
}
// Used by targetDataBegin, targetDataEnd, targetDataUpdate and target.
@@ -365,6 +370,7 @@ DeviceTy::getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast,
void *TargetPointer = NULL;
bool IsNew = false;
+ bool IsPresent = true;
IsHostPtr = false;
IsLast = false;
LookupResult LR = lookupMapping(HDTTMap, HstPtrBegin, Size);
@@ -416,11 +422,18 @@ DeviceTy::getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast,
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
index 197a5489c11d6..865d3a6412f9f 100644
--- a/openmp/libomptarget/src/omptarget.cpp
+++ b/openmp/libomptarget/src/omptarget.cpp
@@ -744,7 +744,8 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
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 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
// 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 @@ static int targetDataContiguous(ident_t *Loc, DeviceTy &Device, void *ArgsBase,
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 @@ static int processDataBefore(ident_t *Loc, int64_t DeviceId, void *HostPtr,
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
index 0000000000000..baa75d21686a7
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/implicit_device_ptr.c
@@ -0,0 +1,26 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+// 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());
+}
More information about the Openmp-commits
mailing list