[Openmp-commits] [openmp] 8a880db - [libomptarget] Make omp_target_is_present checks storage instead of zero length array.

[Ye Luo via Openmp-commits]

Author: Ye Luo
Date: 2022-04-22T17:37:06-05:00
New Revision: 8a880db5192970e80f77130b9e61e548966d96c1

URL: https://github.com/llvm/llvm-project/commit/8a880db5192970e80f77130b9e61e548966d96c1
DIFF: https://github.com/llvm/llvm-project/commit/8a880db5192970e80f77130b9e61e548966d96c1.diff

LOG: [libomptarget] Make omp_target_is_present checks storage instead of zero length array.

Consider checking whether a pointer has been mapped can be achieved via omp_get_mapped_ptr.
omp_target_is_present is more needed to check whether the storage being pointed is mapped.
This restore the old behavior of omp_target_is_present before D123093
Fixes https://github.com/llvm/llvm-project/issues/54899

Reviewed By: jdenny

Differential Revision: https://reviews.llvm.org/D123891

Added: 
    

Modified: 
    openmp/libomptarget/src/api.cpp
    openmp/libomptarget/test/mapping/target_implicit_partial_map.c

Removed: 
    


################################################################################
diff  --git a/openmp/libomptarget/src/api.cpp b/openmp/libomptarget/src/api.cpp
index ba72e2fa135d6..29af0167e1f3d 100644
--- a/openmp/libomptarget/src/api.cpp
+++ b/openmp/libomptarget/src/api.cpp
@@ -108,8 +108,12 @@ EXTERN int omp_target_is_present(const void *ptr, int device_num) {
   DeviceTy &Device = *PM->Devices[device_num];
   bool IsLast; // not used
   bool IsHostPtr;
+  // omp_target_is_present tests whether a host pointer refers to storage that
+  // is mapped to a given device. However, due to the lack of the storage size,
+  // only check 1 byte. Cannot set size 0 which checks whether the pointer (zero
+  // lengh array) is mapped instead of the referred storage.
   TargetPointerResultTy TPR =
-      Device.getTgtPtrBegin(const_cast<void *>(ptr), 0, IsLast,
+      Device.getTgtPtrBegin(const_cast<void *>(ptr), 1, IsLast,
                             /*UpdateRefCount=*/false,
                             /*UseHoldRefCount=*/false, IsHostPtr);
   int rc = (TPR.TargetPointer != NULL);

diff  --git a/openmp/libomptarget/test/mapping/target_implicit_partial_map.c b/openmp/libomptarget/test/mapping/target_implicit_partial_map.c
index 1e9dbb1b0f7bb..8df2a4d8e5857 100644
--- a/openmp/libomptarget/test/mapping/target_implicit_partial_map.c
+++ b/openmp/libomptarget/test/mapping/target_implicit_partial_map.c
@@ -13,6 +13,18 @@ int main() {
 
 #pragma omp target data map(alloc: arr[50:2]) // partially mapped
   {
+    // CHECK: arr[50] must present: 1
+    fprintf(stderr, "arr[50] must present: %d\n",
+            omp_target_is_present(&arr[50], omp_get_default_device()));
+
+    // CHECK: arr[0] should not present: 0
+    fprintf(stderr, "arr[0] should not present: %d\n",
+            omp_target_is_present(&arr[0], omp_get_default_device()));
+
+    // CHECK: arr[49] should not present: 0
+    fprintf(stderr, "arr[49] should not present: %d\n",
+            omp_target_is_present(&arr[49], omp_get_default_device()));
+
 #pragma omp target // would implicitly map with full size but already present
     {
       arr[50] = 5;
@@ -20,8 +32,8 @@ int main() {
     } // must treat as present (dec ref count) even though full size not present
   } // wouldn't delete if previous ref count dec didn't happen
 
-  // CHECK: still present: 0
-  fprintf(stderr, "still present: %d\n",
+  // CHECK: arr[50] still present: 0
+  fprintf(stderr, "arr[50] still present: %d\n",
           omp_target_is_present(&arr[50], omp_get_default_device()));
 
   return 0;