[PATCH v2 7/8] openmp, libgomp: Handle unified shared memory in omp_target_is_accessible

Andrew Stubbs ams@baylibre.com
Fri Jun 28 10:24:48 GMT 2024


From: Marcel Vollweiler <marcel@codesourcery.com>

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.
---
 libgomp/target.c                              |  8 +++++--
 .../target-is-accessible-1.c                  | 22 +++++++++++++------
 .../target-is-accessible-2.c                  | 21 ++++++++++++++++++
 .../target-is-accessible-1.f90                | 20 +++++++++++------
 .../target-is-accessible-2.f90                | 22 +++++++++++++++++++
 5 files changed, 77 insertions(+), 16 deletions(-)
 create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-2.c
 create mode 100644 libgomp/testsuite/libgomp.fortran/target-is-accessible-2.f90

diff --git a/libgomp/target.c b/libgomp/target.c
index 754dea4e031..f0ee2c84197 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -5281,9 +5281,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..24c77232f5d
--- /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_gnu_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_gnu_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..66e5a632961
--- /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_gnu_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_gnu_unified_shared_mem_alloc);
+end program main
-- 
2.41.0



More information about the Gcc-patches mailing list