[gcc/devel/omp/gcc-12] OpenMP, libgomp: Handle unified shared memory in omp_target_is_accessible

[gcc/devel/omp/gcc-12] OpenMP, libgomp: Handle unified shared memory in omp_target_is_accessible

[Marcel Vollweiler]

https://gcc.gnu.org/g:9044b7efb3518de180a5b3168615b7e12d93eea8

commit 9044b7efb3518de180a5b3168615b7e12d93eea8
Author: Marcel Vollweiler <marcel@codesourcery.com>
Date:   Tue Dec 13 12:04:48 2022 +0000

    OpenMP, libgomp: Handle unified shared memory in omp_target_is_accessible
    
    This patch handles Unified Shared Memory (USM) in the OpenMP runtime routine
    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:
---
 libgomp/ChangeLog.omp                              |  8 ++++++++
 libgomp/target.c                                   |  8 ++++++--
 .../libgomp.c-c++-common/target-is-accessible-1.c  | 22 +++++++++++++++-------
 .../libgomp.c-c++-common/target-is-accessible-2.c  | 21 +++++++++++++++++++++
 .../libgomp.fortran/target-is-accessible-1.f90     | 20 +++++++++++++-------
 .../libgomp.fortran/target-is-accessible-2.f90     | 22 ++++++++++++++++++++++
 6 files changed, 85 insertions(+), 16 deletions(-)

diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 32bcc842af8..a0d0271f63b 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,11 @@
+2022-12-13  Marcel Vollweiler  <marcel@codesourcery.com>
+
+	* 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.
+
 2022-12-12  Tobias Burnus  <tobias@codesourcery.com>
 
 	Backported from master:
diff --git a/libgomp/target.c b/libgomp/target.c
index 50709f0677d..2cd8e2a65dd 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -5067,9 +5067,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 2e75c6300ae..e7f9cf27a42 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
@@ -1,3 +1,5 @@
+/* { dg-do run } */
+
 #include <omp.h>
 
 int
@@ -6,7 +8,8 @@ main ()
   int d = omp_get_default_device ();
   int id = omp_get_initial_device ();
   int n = omp_get_num_devices ();
-  void *p;
+  int i = 42;
+  void *p = &i;
 
   if (d < 0 || d >= n)
     d = id;
@@ -26,23 +29,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 00000000000..0917365379b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-2.c
@@ -0,0 +1,21 @@
+/* { dg-do run } */
+/* { dg-require-effective-target omp_usm } */
+
+#include <omp.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 150df6f8a4f..0df43aae095 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
@@ -28,24 +30,28 @@ program main
   if (omp_target_is_accessible (p, c_sizeof (d), n + 1) /= 0) &
     stop 5
 
-  ! 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 6;
 
-    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 7;
 
+    if (shared_mem == 1 .and. omp_target_is_accessible (c_loc (a), 128 * sizeof (a(1)), d) == 0) &
+      stop 8;
+
     do i = 1, 128
-      if (omp_target_is_accessible (c_loc (a(i)), sizeof (a(i)), d) /= shared_mem) &
-        stop 8;
+      if (shared_mem == 1 .and. omp_target_is_accessible (c_loc (a(i)), sizeof (a(i)), d) == 0) &
+        stop 9;
     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 00000000000..624d1efa08e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-is-accessible-2.f90
@@ -0,0 +1,22 @@
+! { dg-do run }
+! { dg-require-effective-target omp_usm }
+
+program main
+  use omp_lib
+  use iso_c_binding
+  implicit none (external, type)
+  integer :: d
+  type(c_ptr) :: p
+
+  !$omp requires unified_shared_memory
+
+  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