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

[ye-luo]

Summary

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

[jdenny]

LGTM based on your discussion so far with the language committee, but please give Ravi some time to comment.

openmp/libomptarget/test/mapping/target_implicit_partial_map.c

19

I suggest checking another boundary case too: &arr[49].

[ye-luo]

Expand the test.

[RaviNarayanaswamy]

LGTM

openmp/libomptarget/src/api.cpp

if (DevicesSize <= (size_t)device_num) {
     DP("Call to omp_target_is_present with invalid device ID, returning "
        "false\n");
     return false;
   }
 
   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);
   // 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)

openmp/libomptarget/test/mapping/target_implicit_partial_map.c

 // RUN: %libomptarget-compile-generic
 // RUN: %libomptarget-run-generic 2>&1 \
 // RUN: | %fcheck-generic
 
 
 // END.
 
 #include <omp.h>
 #include <stdio.h>
 
 int main() {
   int arr[100];
 
 #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;
       arr[51] = 6;
     } // 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;
 }