public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-12] OpenMP, libgomp: Add new runtime routine omp_target_is_accessible.
@ 2022-06-29 14:47 Kwok Yeung
0 siblings, 0 replies; only message in thread
From: Kwok Yeung @ 2022-06-29 14:47 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:caa8134568c639dfe980ebab26a514e7e1fa5199
commit caa8134568c639dfe980ebab26a514e7e1fa5199
Author: Marcel Vollweiler <marcel@codesourcery.com>
Date: Fri May 6 07:28:26 2022 -0700
OpenMP, libgomp: Add new runtime routine omp_target_is_accessible.
gcc/ChangeLog:
* omp-low.cc (omp_runtime_api_call): Added target_is_accessible to
omp_runtime_apis array.
libgomp/ChangeLog:
* libgomp.map: Added omp_target_is_accessible.
* libgomp.texi: Tagged omp_target_is_accessible as supported.
* omp.h.in: Added omp_target_is_accessible.
* omp_lib.f90.in: Added interface for omp_target_is_accessible.
* omp_lib.h.in: Likewise.
* target.c (omp_target_is_accessible): Added implementation of
omp_target_is_accessible.
* testsuite/libgomp.c-c++-common/target-is-accessible-1.c: New test.
* testsuite/libgomp.fortran/target-is-accessible-1.f90: New test.
(cherry picked from commit 4043f53cb4a3541f7a6e4f4132419f78ab7ec4f7)
Diff:
---
gcc/ChangeLog.omp | 8 ++++
gcc/omp-low.cc | 1 +
libgomp/ChangeLog.omp | 15 +++++++
libgomp/libgomp.map | 1 +
libgomp/libgomp.texi | 2 +-
libgomp/omp.h.in | 2 +
libgomp/omp_lib.f90.in | 10 +++++
libgomp/omp_lib.h.in | 11 +++++
libgomp/target.c | 18 ++++++++
.../libgomp.c-c++-common/target-is-accessible-1.c | 47 ++++++++++++++++++++
.../libgomp.fortran/target-is-accessible-1.f90 | 50 ++++++++++++++++++++++
11 files changed, 164 insertions(+), 1 deletion(-)
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index ac8455b6a3e..59ad9943ae0 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,11 @@
+2022-05-06 Marcel Vollweiler <marcel@codesourcery.com>
+
+ Backport from mainline:
+ 2022-05-06 Marcel Vollweiler <marcel@codesourcery.com>
+
+ * omp-low.cc (omp_runtime_api_call): Added target_is_accessible to
+ omp_runtime_apis array.
+
2022-05-05 Sandra Loosemore <sandra@codesourcery.com>
Backport from mainline:
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 6fa7b31f372..bd7515d6b78 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -4488,6 +4488,7 @@ omp_runtime_api_call (const_tree fndecl)
"target_associate_ptr",
"target_disassociate_ptr",
"target_free",
+ "target_is_accessible",
"target_is_present",
"target_memcpy",
"target_memcpy_rect",
diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 327b396ca58..d8a7ac2c8ef 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,18 @@
+2022-05-06 Marcel Vollweiler <marcel@codesourcery.com>
+
+ Backport from mainline:
+ 2022-05-06 Marcel Vollweiler <marcel@codesourcery.com>
+
+ * libgomp.map: Added omp_target_is_accessible.
+ * libgomp.texi: Tagged omp_target_is_accessible as supported.
+ * omp.h.in: Added omp_target_is_accessible.
+ * omp_lib.f90.in: Added interface for omp_target_is_accessible.
+ * omp_lib.h.in: Likewise.
+ * target.c (omp_target_is_accessible): Added implementation of
+ omp_target_is_accessible.
+ * testsuite/libgomp.c-c++-common/target-is-accessible-1.c: New test.
+ * testsuite/libgomp.fortran/target-is-accessible-1.f90: New test.
+
2022-05-05 Sandra Loosemore <sandra@codesourcery.com>
Backport from mainline:
diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map
index fc14b8d172b..4423f9c0965 100644
--- a/libgomp/libgomp.map
+++ b/libgomp/libgomp.map
@@ -229,6 +229,7 @@ OMP_5.1 {
OMP_5.1.1 {
global:
omp_get_mapped_ptr;
+ omp_target_is_accessible;
} OMP_5.1;
GOMP_1.0 {
diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
index 317711f3f73..ca959b644d4 100644
--- a/libgomp/libgomp.texi
+++ b/libgomp/libgomp.texi
@@ -311,7 +311,7 @@ The OpenMP 4.5 specification is fully supported.
@item @code{omp_set_num_teams}, @code{omp_set_teams_thread_limit},
@code{omp_get_max_teams}, @code{omp_get_teams_thread_limit} runtime
routines @tab Y @tab
-@item @code{omp_target_is_accessible} runtime routine @tab N @tab
+@item @code{omp_target_is_accessible} runtime routine @tab Y @tab
@item @code{omp_target_memcpy_async} and @code{omp_target_memcpy_rect_async}
runtime routines @tab N @tab
@item @code{omp_get_mapped_ptr} runtime routine @tab Y @tab
diff --git a/libgomp/omp.h.in b/libgomp/omp.h.in
index e811911c99b..5640ad2a70f 100644
--- a/libgomp/omp.h.in
+++ b/libgomp/omp.h.in
@@ -288,6 +288,8 @@ extern int omp_target_associate_ptr (const void *, const void *, __SIZE_TYPE__,
__SIZE_TYPE__, int) __GOMP_NOTHROW;
extern int omp_target_disassociate_ptr (const void *, int) __GOMP_NOTHROW;
extern void *omp_get_mapped_ptr (const void *, int) __GOMP_NOTHROW;
+extern int omp_target_is_accessible (const void *, __SIZE_TYPE__, int)
+ __GOMP_NOTHROW;
extern void omp_set_affinity_format (const char *) __GOMP_NOTHROW;
extern __SIZE_TYPE__ omp_get_affinity_format (char *, __SIZE_TYPE__)
diff --git a/libgomp/omp_lib.f90.in b/libgomp/omp_lib.f90.in
index d6a251799ac..3b2461819eb 100644
--- a/libgomp/omp_lib.f90.in
+++ b/libgomp/omp_lib.f90.in
@@ -854,6 +854,16 @@
end function omp_get_mapped_ptr
end interface
+ interface
+ function omp_target_is_accessible (ptr, size, device_num) bind(c)
+ use, intrinsic :: iso_c_binding, only : c_ptr, c_size_t, c_int
+ integer(c_int) :: omp_target_is_accessible
+ type(c_ptr), value :: ptr
+ integer(c_size_t), value :: size
+ integer(c_int), value :: device_num
+ end function omp_target_is_accessible
+ end interface
+
#if _OPENMP >= 201811
!GCC$ ATTRIBUTES DEPRECATED :: omp_get_nested, omp_set_nested
#endif
diff --git a/libgomp/omp_lib.h.in b/libgomp/omp_lib.h.in
index 0f48510d7ff..28554331c95 100644
--- a/libgomp/omp_lib.h.in
+++ b/libgomp/omp_lib.h.in
@@ -425,3 +425,14 @@
integer(c_int), value :: device_num
end function omp_get_mapped_ptr
end interface
+
+ interface
+ function omp_target_is_accessible (ptr, size, device_num) &
+ & bind(c)
+ use, intrinsic :: iso_c_binding, only : c_ptr, c_size_t, c_int
+ integer(c_int) :: omp_target_is_accessible
+ type(c_ptr), value :: ptr
+ integer(c_size_t), value :: size
+ integer(c_int), value :: device_num
+ end function omp_target_is_accessible
+ end interface
diff --git a/libgomp/target.c b/libgomp/target.c
index b6298f80962..9a3b44f6e1b 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -3997,6 +3997,24 @@ omp_get_mapped_ptr (const void *ptr, int device_num)
return ret;
}
+int
+omp_target_is_accessible (const void *ptr, size_t size, int device_num)
+{
+ if (device_num < 0 || device_num > gomp_get_num_devices ())
+ return false;
+
+ if (device_num == gomp_get_num_devices ())
+ return true;
+
+ struct gomp_device_descr *devicep = resolve_device (device_num);
+ if (devicep == NULL)
+ return false;
+
+ /* TODO: Unified shared memory must be handled when available. */
+
+ return devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM;
+}
+
int
omp_pause_resource (omp_pause_resource_t kind, int device_num)
{
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
new file mode 100644
index 00000000000..7c2cf622960
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-is-accessible-1.c
@@ -0,0 +1,47 @@
+#include <omp.h>
+
+int
+main ()
+{
+ int d = omp_get_default_device ();
+ int id = omp_get_initial_device ();
+ int n = omp_get_num_devices ();
+ void *p;
+
+ if (d < 0 || d >= n)
+ d = id;
+
+ if (!omp_target_is_accessible (p, sizeof (int), n))
+ __builtin_abort ();
+
+ if (!omp_target_is_accessible (p, sizeof (int), id))
+ __builtin_abort ();
+
+ if (omp_target_is_accessible (p, sizeof (int), -1))
+ __builtin_abort ();
+
+ 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++)
+ {
+ 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)
+ __builtin_abort ();
+
+ if (omp_target_is_accessible (a, 128 * sizeof (int), d) != shared_mem)
+ __builtin_abort ();
+
+ for (int i = 0; i < 128; i++)
+ if (omp_target_is_accessible (&a[i], sizeof (int), d) != shared_mem)
+ __builtin_abort ();
+ }
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.fortran/target-is-accessible-1.f90 b/libgomp/testsuite/libgomp.fortran/target-is-accessible-1.f90
new file mode 100644
index 00000000000..26118553f2d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-is-accessible-1.f90
@@ -0,0 +1,50 @@
+program main
+ use omp_lib
+ use iso_c_binding
+ implicit none (external, type)
+ integer :: d, id, n, shared_mem, i
+ integer, target :: a(1:128)
+ type(c_ptr) :: p
+
+ d = omp_get_default_device ()
+ id = omp_get_initial_device ()
+ n = omp_get_num_devices ()
+
+ if (d < 0 .or. d >= n) &
+ d = id
+
+ if (omp_target_is_accessible (p, c_sizeof (d), n) /= 1) &
+ stop 1
+
+ if (omp_target_is_accessible (p, c_sizeof (d), id) /= 1) &
+ stop 2
+
+ if (omp_target_is_accessible (p, c_sizeof (d), -1) /= 0) &
+ stop 3
+
+ 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 = 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) &
+ stop 5;
+
+ if (omp_target_is_accessible (c_loc (a), 128 * sizeof (a(1)), d) /= shared_mem) &
+ stop 6;
+
+ do i = 1, 128
+ if (omp_target_is_accessible (c_loc (a(i)), sizeof (a(i)), d) /= shared_mem) &
+ stop 7;
+ end do
+
+ end do
+
+end program main
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2022-06-29 14:47 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2022-06-29 14:47 [gcc/devel/omp/gcc-12] OpenMP, libgomp: Add new runtime routine omp_target_is_accessible Kwok Yeung
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).