OpenMP, libgomp: Handle unified shared memory in omp_target_is_accessible.

libgomp/ChangeLog:

	* target.c (omp_target_is_accessible): Handle unified shared memory.
	* testsuite/libgomp.c-c++-common/target-is-accessible-1.c: Updated.
	* testsuite/libgomp.fortran/target-is-accessible-1.f90: Updated.
	* testsuite/libgomp.c-c++-common/target-is-accessible-2.c: New test.
	* testsuite/libgomp.fortran/target-is-accessible-2.f90: New test.

diff --git a/libgomp/target.c b/libgomp/target.c
index 74a031f..e6d00c5 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -3909,9 +3909,13 @@ omp_target_is_accessible (const void *ptr, size_t size, int device_num)
   if (devicep == NULL)
     return false;
 
-  /* TODO: Unified shared memory must be handled when available.  */
+  if (devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+    return true;
 
-  return devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM;
+  if (devicep->is_usm_ptr_func && devicep->is_usm_ptr_func ((void *) ptr))
+    return true;
+
+  return false;
 }
 
 int
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-1.c
index 7c2cf62..e3f494b 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-1.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-1.c
@@ -23,23 +23,28 @@ main ()
   if (omp_target_is_accessible (p, sizeof (int), n + 1))
     __builtin_abort ();
 
-  /* Currently, a host pointer is accessible if the device supports shared
-     memory or omp_target_is_accessible is executed on the host. This
-     test case must be adapted when unified shared memory is avialable.  */
   int a[128];
   for (int d = 0; d <= omp_get_num_devices (); d++)
     {
+      /* SHARED_MEM is 1 if and only if host and device share the same memory.
+	 OMP_TARGET_IS_ACCESSIBLE should not return 0 for shared memory.  */
       int shared_mem = 0;
       #pragma omp target map (alloc: shared_mem) device (d)
 	shared_mem = 1;
-      if (omp_target_is_accessible (p, sizeof (int), d) != shared_mem)
+
+      if (shared_mem && !omp_target_is_accessible (p, sizeof (int), d))
+	__builtin_abort ();
+
+      /* USM is disabled by default.  Hence OMP_TARGET_IS_ACCESSIBLE should
+	 return 0 if shared_mem is false.  */
+      if (!shared_mem && omp_target_is_accessible (p, sizeof (int), d))
 	__builtin_abort ();
 
-      if (omp_target_is_accessible (a, 128 * sizeof (int), d) != shared_mem)
+      if (shared_mem && !omp_target_is_accessible (a, 128 * sizeof (int), d))
 	__builtin_abort ();
 
       for (int i = 0; i < 128; i++)
-	if (omp_target_is_accessible (&a[i], sizeof (int), d) != shared_mem)
+	if (shared_mem && !omp_target_is_accessible (&a[i], sizeof (int), d))
 	  __builtin_abort ();
     }
 
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-2.c
new file mode 100644
index 0000000..24af51f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-2.c
@@ -0,0 +1,22 @@
+/* { dg-do run } */
+/* { dg-skip-if "USM is only implemented for nvptx." { ! offload_target_nvptx } } */
+
+#include <omp.h>
+#include <stdint.h>
+
+#pragma omp requires unified_shared_memory
+
+int
+main ()
+{
+  int *a = (int *) omp_alloc (sizeof(int), ompx_unified_shared_mem_alloc);
+  if (!a)
+    __builtin_abort ();
+
+  for (int d = 0; d <= omp_get_num_devices (); d++)
+    if (!omp_target_is_accessible (a, sizeof (int), d))
+      __builtin_abort ();
+
+  omp_free(a, ompx_unified_shared_mem_alloc);
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.fortran/target-is-accessible-1.f90 b/libgomp/testsuite/libgomp.fortran/target-is-accessible-1.f90
index 2611855..015f74a 100644
--- a/libgomp/testsuite/libgomp.fortran/target-is-accessible-1.f90
+++ b/libgomp/testsuite/libgomp.fortran/target-is-accessible-1.f90
@@ -1,3 +1,5 @@
+! { dg-do run }
+
 program main
   use omp_lib
   use iso_c_binding
@@ -25,24 +27,28 @@ program main
   if (omp_target_is_accessible (p, c_sizeof (d), n + 1) /= 0) &
     stop 4
 
-  ! Currently, a host pointer is accessible if the device supports shared
-  ! memory or omp_target_is_accessible is executed on the host. This
-  ! test case must be adapted when unified shared memory is avialable.
   do d = 0, omp_get_num_devices ()
+    ! SHARED_MEM is 1 if and only if host and device share the same memory.
+    ! OMP_TARGET_IS_ACCESSIBLE should not return 0 for shared memory.
     shared_mem = 0;
     !$omp target map (alloc: shared_mem) device (d)
       shared_mem = 1;
     !$omp end target
 
-    if (omp_target_is_accessible (p, c_sizeof (d), d) /= shared_mem) &
+    if (shared_mem == 1 .and. omp_target_is_accessible (p, c_sizeof (d), d) == 0) &
       stop 5;
 
-    if (omp_target_is_accessible (c_loc (a), 128 * sizeof (a(1)), d) /= shared_mem) &
+    ! USM is disabled by default.  Hence OMP_TARGET_IS_ACCESSIBLE should
+    ! return 0 if shared_mem is false.
+    if (shared_mem == 0 .and. omp_target_is_accessible (p, c_sizeof (d), d) /= 0) &
       stop 6;
 
+    if (shared_mem == 1 .and. omp_target_is_accessible (c_loc (a), 128 * sizeof (a(1)), d) == 0) &
+      stop 7;
+
     do i = 1, 128
-      if (omp_target_is_accessible (c_loc (a(i)), sizeof (a(i)), d) /= shared_mem) &
-        stop 7;
+      if (shared_mem == 1 .and. omp_target_is_accessible (c_loc (a(i)), sizeof (a(i)), d) == 0) &
+        stop 8;
     end do
 
   end do
diff --git a/libgomp/testsuite/libgomp.fortran/target-is-accessible-2.f90 b/libgomp/testsuite/libgomp.fortran/target-is-accessible-2.f90
new file mode 100644
index 0000000..5c08564
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-is-accessible-2.f90
@@ -0,0 +1,20 @@
+! { dg-do run }
+! { dg-skip-if "USM is only implemented for nvptx." { ! offload_target_nvptx } }
+
+program main
+  use omp_lib
+  use iso_c_binding
+  implicit none (external, type)
+  integer :: d
+  type(c_ptr) :: p
+
+  p = omp_alloc (sizeof(d), ompx_unified_shared_mem_alloc)
+  if (.not. c_associated (p)) stop 1
+
+  do d = 0, omp_get_num_devices ()
+    if (omp_target_is_accessible (p, c_sizeof (d), d) == 0) &
+      stop 2;
+  end do
+
+  call omp_free (p, ompx_unified_shared_mem_alloc);
+end program main