* [PATCH] OpenMP, libgomp: Handle unified shared memory in omp_target_is_accessible.
@ 2022-05-06 11:19 Marcel Vollweiler
0 siblings, 0 replies; only message in thread
From: Marcel Vollweiler @ 2022-05-06 11:19 UTC (permalink / raw)
To: gcc-patches; +Cc: Jakub Jelinek, fortran
[-- Attachment #1: Type: text/plain, Size: 649 bytes --]
Hi,
This is a follow up patch of the patch that adds the OpenMP runtime routine
omp_target_is_accessible:
https://gcc.gnu.org/pipermail/gcc-patches/2022-March/591601.html
It considers now also unified shared memory (usm) that was submitted recently
(but not yet approved/committed):
https://gcc.gnu.org/pipermail/gcc-patches/2022-March/591349.html
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
[-- Attachment #2: target-is-accessible-with-usm-patch.diff --]
[-- Type: text/plain, Size: 6294 bytes --]
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
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2022-05-06 11:20 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2022-05-06 11:19 [PATCH] OpenMP, libgomp: Handle unified shared memory in omp_target_is_accessible Marcel Vollweiler
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).