diff mbox series

[OG12,committed] OpenMP, libgomp: Handle unified shared memory in omp_target_is_accessible.

Message ID 0ff30ea1-ee5f-7f10-dcbc-bea85e2bfa81@codesourcery.com
State Unresolved
Headers
Series [OG12,committed] OpenMP, libgomp: Handle unified shared memory in omp_target_is_accessible. |

Checks

Context Check Description
snail/gcc-patch-check warning Git am fail log

Commit Message

Marcel Vollweiler Dec. 13, 2022, 4:12 p.m. UTC 
  This patch handles Unified Shared Memory (USM) in the OpenMP runtime routine
omp_target_is_accessible implementation.

A previous patch was submitted some months ago
(https://gcc.gnu.org/pipermail/gcc-patches/2022-May/594187.html) but not yet
reviewed due to dependencies on the Unified Shared Memory implementation.
Although USM is not yet in mainline, the corresponding patches were already
committed to OG12. I rebased, updated, and committed my patch to OG12
(devel/omp/gcc-12 branch).

I tested the patch with nvptx offloading (x86_64-linux and PowerPC) without
regressions. Since USM is not supported for all gcn targets, I tested gcn with
offloading for x86_64-linux on various targets (gfx90a, gfx908, gfx906, gfx803)
- also without regressions.

Marcel
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
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 mbox series

Patch

diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 32bcc84..a0d0271 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 50709f0..2cd8e2a 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 2e75c63..e7f9cf2 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 0000000..0917365
--- /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 150df6f..0df43aae 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 0000000..624d1ef
--- /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