public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-12] OpenMP, libgomp: Add new runtime routine omp_get_mapped_ptr.
@ 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:28c49e0f617b50b80e974582479b7fe077988b6d
commit 28c49e0f617b50b80e974582479b7fe077988b6d
Author: Marcel Vollweiler <marcel@codesourcery.com>
Date: Mon May 2 23:56:44 2022 -0700
OpenMP, libgomp: Add new runtime routine omp_get_mapped_ptr.
This patch adds the OpenMP runtime routine "omp_get_mapped_ptr" which was
introduced in OpenMP 5.1.
gcc/ChangeLog:
* omp-low.cc (omp_runtime_api_call): Added get_mapped_ptr to
omp_runtime_apis array.
libgomp/ChangeLog:
* libgomp.map: Added omp_get_mapped_ptr.
* libgomp.texi: Tagged omp_get_mapped_ptr as supported.
* omp.h.in: Added omp_get_mapped_ptr.
* omp_lib.f90.in: Added interface for omp_get_mapped_ptr.
* omp_lib.h.in: Likewise.
* target.c (omp_get_mapped_ptr): Added implementation of
omp_get_mapped_ptr.
* testsuite/libgomp.c-c++-common/get-mapped-ptr-1.c: New test.
* testsuite/libgomp.c-c++-common/get-mapped-ptr-2.c: New test.
* testsuite/libgomp.c-c++-common/get-mapped-ptr-3.c: New test.
* testsuite/libgomp.c-c++-common/get-mapped-ptr-4.c: New test.
* testsuite/libgomp.fortran/get-mapped-ptr-1.f90: New test.
* testsuite/libgomp.fortran/get-mapped-ptr-2.f90: New test.
* testsuite/libgomp.fortran/get-mapped-ptr-3.f90: New test.
* testsuite/libgomp.fortran/get-mapped-ptr-4.f90: New test.
(cherry picked from commit 941cdc8b6d29f9fe494fdd244e96a5e5aa08ba32)
Diff:
---
gcc/ChangeLog.omp | 8 +
gcc/omp-low.cc | 1 +
libgomp/ChangeLog.omp | 21 +++
libgomp/libgomp.map | 5 +
libgomp/libgomp.texi | 2 +-
libgomp/omp.h.in | 1 +
libgomp/omp_lib.f90.in | 9 ++
libgomp/omp_lib.h.in | 9 ++
libgomp/target.c | 38 +++++
.../libgomp.c-c++-common/get-mapped-ptr-1.c | 41 +++++
.../libgomp.c-c++-common/get-mapped-ptr-2.c | 106 +++++++++++++
.../libgomp.c-c++-common/get-mapped-ptr-3.c | 51 ++++++
.../libgomp.c-c++-common/get-mapped-ptr-4.c | 49 ++++++
.../testsuite/libgomp.fortran/get-mapped-ptr-1.f90 | 43 +++++
.../testsuite/libgomp.fortran/get-mapped-ptr-2.f90 | 175 +++++++++++++++++++++
.../testsuite/libgomp.fortran/get-mapped-ptr-3.f90 | 48 ++++++
.../testsuite/libgomp.fortran/get-mapped-ptr-4.f90 | 84 ++++++++++
17 files changed, 690 insertions(+), 1 deletion(-)
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index 12441742a61..916f4a896b5 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,11 @@
+2022-05-02 Marcel Vollweiler <marcel@codesourcery.com>
+
+ Backport from mainline:
+ 2022-05-02 Marcel Vollweiler <marcel@codesourcery.com>
+
+ * omp-low.cc (omp_runtime_api_call): Added get_mapped_ptr to
+ omp_runtime_apis array.
+
2022-06-17 Chung-Lin Tang <cltang@codesourcery.com>
Backport from mainline:
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 2838f520390..6fa7b31f372 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -4482,6 +4482,7 @@ omp_runtime_api_call (const_tree fndecl)
"alloc",
"calloc",
"free",
+ "get_mapped_ptr",
"realloc",
"target_alloc",
"target_associate_ptr",
diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 5bbd997dcce..5c53648f6be 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,24 @@
+2022-05-02 Marcel Vollweiler <marcel@codesourcery.com>
+
+ Backport from mainline:
+ 2022-05-02 Marcel Vollweiler <marcel@codesourcery.com>
+
+ * libgomp.map: Added omp_get_mapped_ptr.
+ * libgomp.texi: Tagged omp_get_mapped_ptr as supported.
+ * omp.h.in: Added omp_get_mapped_ptr.
+ * omp_lib.f90.in: Added interface for omp_get_mapped_ptr.
+ * omp_lib.h.in: Likewise.
+ * target.c (omp_get_mapped_ptr): Added implementation of
+ omp_get_mapped_ptr.
+ * testsuite/libgomp.c-c++-common/get-mapped-ptr-1.c: New test.
+ * testsuite/libgomp.c-c++-common/get-mapped-ptr-2.c: New test.
+ * testsuite/libgomp.c-c++-common/get-mapped-ptr-3.c: New test.
+ * testsuite/libgomp.c-c++-common/get-mapped-ptr-4.c: New test.
+ * testsuite/libgomp.fortran/get-mapped-ptr-1.f90: New test.
+ * testsuite/libgomp.fortran/get-mapped-ptr-2.f90: New test.
+ * testsuite/libgomp.fortran/get-mapped-ptr-3.f90: New test.
+ * testsuite/libgomp.fortran/get-mapped-ptr-4.f90: New test.
+
2022-02-24 Andrew Stubbs <ams@codesourcery.com>
Backport from mainline:
diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map
index f714900e24d..fc14b8d172b 100644
--- a/libgomp/libgomp.map
+++ b/libgomp/libgomp.map
@@ -226,6 +226,11 @@ OMP_5.1 {
omp_get_teams_thread_limit_;
} OMP_5.0.2;
+OMP_5.1.1 {
+ global:
+ omp_get_mapped_ptr;
+} OMP_5.1;
+
GOMP_1.0 {
global:
GOMP_atomic_end;
diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
index cfd8ca59cd2..2d4e113f1e3 100644
--- a/libgomp/libgomp.texi
+++ b/libgomp/libgomp.texi
@@ -314,7 +314,7 @@ The OpenMP 4.5 specification is fully supported.
@item @code{omp_target_is_accessible} runtime routine @tab N @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 N @tab
+@item @code{omp_get_mapped_ptr} runtime routine @tab Y @tab
@item @code{omp_calloc}, @code{omp_realloc}, @code{omp_aligned_alloc} and
@code{omp_aligned_calloc} runtime routines @tab Y @tab
@item @code{omp_alloctrait_key_t} enum: @code{omp_atv_serialized} added,
diff --git a/libgomp/omp.h.in b/libgomp/omp.h.in
index 4ec4475306b..e811911c99b 100644
--- a/libgomp/omp.h.in
+++ b/libgomp/omp.h.in
@@ -287,6 +287,7 @@ extern int omp_target_memcpy_rect (void *, const void *, __SIZE_TYPE__, int,
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 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 e1c32aa78d2..d6a251799ac 100644
--- a/libgomp/omp_lib.f90.in
+++ b/libgomp/omp_lib.f90.in
@@ -845,6 +845,15 @@
end function omp_target_disassociate_ptr
end interface
+ interface
+ function omp_get_mapped_ptr (ptr, device_num) bind(c)
+ use, intrinsic :: iso_c_binding, only : c_ptr, c_int
+ type(c_ptr) :: omp_get_mapped_ptr
+ type(c_ptr), value :: ptr
+ integer(c_int), value :: device_num
+ end function omp_get_mapped_ptr
+ 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 ff857a479df..0f48510d7ff 100644
--- a/libgomp/omp_lib.h.in
+++ b/libgomp/omp_lib.h.in
@@ -416,3 +416,12 @@
integer(c_int), value :: device_num
end function omp_target_disassociate_ptr
end interface
+
+ interface
+ function omp_get_mapped_ptr (ptr, device_num) bind(c)
+ use, intrinsic :: iso_c_binding, only : c_ptr, c_int
+ type(c_ptr) :: omp_get_mapped_ptr
+ type(c_ptr), value :: ptr
+ integer(c_int), value :: device_num
+ end function omp_get_mapped_ptr
+ end interface
diff --git a/libgomp/target.c b/libgomp/target.c
index 93e904725b8..b6298f80962 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -3959,6 +3959,44 @@ omp_target_disassociate_ptr (const void *ptr, int device_num)
return ret;
}
+void *
+omp_get_mapped_ptr (const void *ptr, int device_num)
+{
+ if (device_num < 0 || device_num > gomp_get_num_devices ())
+ return NULL;
+
+ if (device_num == omp_get_initial_device ())
+ return (void *) ptr;
+
+ struct gomp_device_descr *devicep = resolve_device (device_num);
+ if (devicep == NULL)
+ return NULL;
+
+ if (!(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
+ || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+ return (void *) ptr;
+
+ gomp_mutex_lock (&devicep->lock);
+
+ struct splay_tree_s *mem_map = &devicep->mem_map;
+ struct splay_tree_key_s cur_node;
+ void *ret = NULL;
+
+ cur_node.host_start = (uintptr_t) ptr;
+ cur_node.host_end = cur_node.host_start;
+ splay_tree_key n = gomp_map_0len_lookup (mem_map, &cur_node);
+
+ if (n)
+ {
+ uintptr_t offset = cur_node.host_start - n->host_start;
+ ret = (void *) (n->tgt->tgt_start + n->tgt_offset + offset);
+ }
+
+ gomp_mutex_unlock (&devicep->lock);
+
+ return ret;
+}
+
int
omp_pause_resource (omp_pause_resource_t kind, int device_num)
{
diff --git a/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-1.c b/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-1.c
new file mode 100644
index 00000000000..97a60ca9541
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-1.c
@@ -0,0 +1,41 @@
+#include <omp.h>
+#include <stdlib.h>
+
+int
+main ()
+{
+ int d = omp_get_default_device ();
+ int id = omp_get_initial_device ();
+ void *p , *q;
+
+ if (d < 0 || d >= omp_get_num_devices ())
+ d = id;
+
+ p = omp_target_alloc (sizeof (int), d);
+ if (p == NULL)
+ return 0;
+
+ if (omp_target_associate_ptr (q, p, sizeof (int), 0, d) != 0)
+ return 0;
+
+ if (omp_get_mapped_ptr (q, -1) != NULL)
+ abort ();
+
+ if (omp_get_mapped_ptr (q, omp_get_num_devices () + 1) != NULL)
+ abort ();
+
+ if (omp_get_mapped_ptr (q, id) != q)
+ abort ();
+
+ if (omp_get_mapped_ptr (q, d) != p)
+ abort ();
+
+ if (omp_target_disassociate_ptr (q, d) != 0)
+ abort ();
+
+ if (omp_get_mapped_ptr (q, d) != NULL)
+ abort ();
+
+ omp_target_free (p, d);
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-2.c b/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-2.c
new file mode 100644
index 00000000000..194dade8ac5
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-2.c
@@ -0,0 +1,106 @@
+#include <omp.h>
+#include <stdlib.h>
+#include <stdint.h>
+
+int
+main ()
+{
+ int d = omp_get_default_device ();
+ int id = omp_get_initial_device ();
+ int a = 42;
+ int b[] = { 24, 42 };
+ int c[] = { 47, 11 };
+ int e[128];
+ int *q = &a;
+ void *p1 = NULL, *p2 = NULL, *p3 = NULL;
+ void *devptrs[128];
+
+ if (d < 0 || d >= omp_get_num_devices ())
+ d = id;
+
+ for (int i = 0; i < 128; i++)
+ e[i] = i;
+
+ #pragma omp target data map(alloc: a, b, c[1], e[32:64]) device(d)
+ {
+ #pragma omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c[1], e[32:64]) device(d)
+ {
+ p1 = &a;
+ p2 = &b;
+ p3 = &c[1];
+ for (int i = 32; i < 96; i++)
+ devptrs[i] = &e[i];
+ }
+
+ if (omp_get_mapped_ptr (&a, d) != (d == id ? &a : p1)
+ || omp_get_mapped_ptr (q, d) != (d == id ? q : p1)
+ || omp_get_mapped_ptr (b, d) != (d == id ? b : p2)
+ || omp_get_mapped_ptr (&b[0], d) != (d == id ? &b[0] : p2)
+ || omp_get_mapped_ptr (&c[1], d) != (d == id ? &c[1] : p3)
+ || omp_get_mapped_ptr (&c[0], d) != (d == id ? &c[0] : NULL))
+ abort ();
+
+ for (int i = 0; i < 32; i++)
+ if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : NULL))
+ abort ();
+ for (int i = 32; i < 96; i++)
+ if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : devptrs[i]))
+ abort ();
+ for (int i = 96; i < 128; i++)
+ if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : NULL))
+ abort ();
+ }
+
+ if (omp_get_mapped_ptr (&a, d) != (d == id ? &a : NULL)
+ || omp_get_mapped_ptr (q, d) != (d == id ? q : NULL)
+ || omp_get_mapped_ptr (b, d) != (d == id ? b : NULL)
+ || omp_get_mapped_ptr (&b[0], d) != (d == id ? &b[0] : NULL)
+ || omp_get_mapped_ptr (&c[1], d) != (d == id ? &c[1] : NULL)
+ || omp_get_mapped_ptr (&c[0], d) != (d == id ? &c[0] : NULL))
+ abort ();
+ for (int i = 0; i < 128; i++)
+ if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : NULL))
+ abort ();
+
+ #pragma omp target enter data map (alloc: a, b, c[1], e[32:64]) device (d)
+ #pragma omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c[1], e[32:64]) device(d)
+ {
+ p1 = &a;
+ p2 = &b;
+ p3 = &c[1];
+ for (int i = 32; i < 96; i++)
+ devptrs[i] = &e[i];
+ }
+
+ if (omp_get_mapped_ptr (&a, d) != (d == id ? &a : p1)
+ || omp_get_mapped_ptr (q, d) != (d == id ? q : p1)
+ || omp_get_mapped_ptr (b, d) != (d == id ? b : p2)
+ || omp_get_mapped_ptr (&b[0], d) != (d == id ? &b[0] : p2)
+ || omp_get_mapped_ptr (&c[1], d) != (d == id ? &c[1] : p3)
+ || omp_get_mapped_ptr (&c[0], d) != (d == id ? &c[0] : NULL))
+ abort ();
+ for (int i = 0; i < 32; i++)
+ if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : NULL))
+ abort ();
+ for (int i = 32; i < 96; i++)
+ if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : devptrs[i]))
+ abort ();
+ for (int i = 96; i < 128; i++)
+ if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : NULL))
+ abort ();
+
+ #pragma omp target exit data map (delete: a, b, c[1], e[32:64]) device (d)
+
+ if (omp_get_mapped_ptr (&a, d) != (d == id ? &a : NULL)
+ || omp_get_mapped_ptr (q, d) != (d == id ? q : NULL)
+ || omp_get_mapped_ptr (b, d) != (d == id ? b : NULL)
+ || omp_get_mapped_ptr (&b[0], d) != (d == id ? &b[0] : NULL)
+ || omp_get_mapped_ptr (&c[1], d) != (d == id ? &c[1] : NULL)
+ || omp_get_mapped_ptr (&c[0], d) != (d == id ? &c[0] : NULL))
+ abort ();
+ for (int i = 0; i < 128; i++)
+ if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : NULL))
+ abort ();
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-3.c b/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-3.c
new file mode 100644
index 00000000000..747ef75c752
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-3.c
@@ -0,0 +1,51 @@
+#include <omp.h>
+#include <stdlib.h>
+
+int
+main ()
+{
+ int d = omp_get_default_device ();
+ int id = omp_get_initial_device ();
+ int a[0];
+ int b[] = { 24, 42 };
+ void *p1 = NULL, *p2 = NULL;
+
+ if (d < 0 || d >= omp_get_num_devices ())
+ d = id;
+
+ void *p = omp_target_alloc (sizeof (int), d);
+ if (p == NULL)
+ return 0;
+
+ if (omp_target_associate_ptr (a, p, sizeof (int), 0, d) != 0)
+ return 0;
+
+ if (omp_get_mapped_ptr (a, d) != (d == id ? a : p))
+ abort ();
+
+ if (omp_target_disassociate_ptr (a, d) != 0)
+ abort ();
+
+ if (omp_get_mapped_ptr (a, d) != (d == id ? a : NULL))
+ abort ();
+
+ #pragma omp target data map(alloc: a, b[1:0]) device(d)
+ {
+ #pragma omp target map(from: p1, p2) map(alloc: a, b[1:0]) device(d)
+ {
+ p1 = &a;
+ p2 = &b[1];
+ }
+
+ /* This is probably expected to be p1/p2 instead of NULL. Zero-length arrays
+ as list items of the map clause are currently not inserted into the mem
+ map ?! However by returning NULL, omp_get_mapped_ptr is consistent with
+ omp_target_is_present. */
+ if (omp_get_mapped_ptr (a, d) != NULL
+ || omp_get_mapped_ptr (&b[1], d) != NULL)
+ abort ();
+ }
+
+ omp_target_free (p, d);
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-4.c b/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-4.c
new file mode 100644
index 00000000000..6f4bd625d48
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-4.c
@@ -0,0 +1,49 @@
+#include <omp.h>
+#include <stdlib.h>
+
+int
+main ()
+{
+ int d = omp_get_default_device ();
+ int id = omp_get_initial_device ();
+ struct s_t { int m1; char m2; } s;
+ void *p1 = NULL, *p2 = NULL;
+
+ if (d < 0 || d >= omp_get_num_devices ())
+ d = id;
+
+ #pragma omp target data map(alloc: s, s.m2) device(d)
+ {
+ #pragma omp target map(from: p1, p2) map(alloc: s, s.m2) device(d)
+ {
+ p1 = &s;
+ p2 = &s.m2;
+ }
+ if (omp_get_mapped_ptr (&s, d) != (d == id ? &s : p1)
+ || omp_get_mapped_ptr (&s.m2, d) != (d == id ? &s.m2 : p2))
+ abort ();
+ }
+
+ if (omp_get_mapped_ptr (&s, d) != (d == id ? &s : NULL)
+ || omp_get_mapped_ptr (&s.m2, d) != (d == id ? &s.m2 : NULL))
+ abort ();
+
+ #pragma omp target enter data map(alloc: s, s.m2) device (d)
+ #pragma omp target map(from: p1, p2) map(alloc: s, s.m2) device(d)
+ {
+ p1 = &s;
+ p2 = &s.m2;
+ }
+
+ if (omp_get_mapped_ptr (&s, d) != (d == id ? &s : p1)
+ || omp_get_mapped_ptr (&s.m2, d) != (d == id ? &s.m2 : p2))
+ abort ();
+
+ #pragma omp target exit data map (delete: s, s.m2) device (d)
+
+ if (omp_get_mapped_ptr (&s, d) != (d == id ? &s : NULL)
+ || omp_get_mapped_ptr (&s.m2, d) != (d == id ? &s.m2 : NULL))
+ abort ();
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-1.f90 b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-1.f90
new file mode 100644
index 00000000000..de05179ce9f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-1.f90
@@ -0,0 +1,43 @@
+program main
+ use omp_lib
+ use iso_c_binding
+ implicit none (external, type)
+ integer :: d, id
+ type(c_ptr) :: p
+ integer, target :: q
+
+ d = omp_get_default_device ()
+ id = omp_get_initial_device ()
+
+ if (d < 0 .or. d >= omp_get_num_devices ()) &
+ d = id
+
+ p = omp_target_alloc (c_sizeof (q), d)
+ if (.not. c_associated (p)) &
+ stop 0 ! okay
+
+ if (omp_target_associate_ptr (c_loc (q), p, c_sizeof (q), &
+ 0_c_size_t, d) == 0) then
+
+ if(c_associated (omp_get_mapped_ptr (c_loc (q), -1))) &
+ stop 1
+
+ if(c_associated (omp_get_mapped_ptr (c_loc (q), &
+ omp_get_num_devices () + 1))) &
+ stop 2
+
+ if(.not. c_associated (omp_get_mapped_ptr (c_loc (q), id), c_loc (q))) &
+ stop 3
+
+ if(.not. c_associated (omp_get_mapped_ptr (c_loc (q), d), p)) &
+ stop 4
+
+ if (omp_target_disassociate_ptr (c_loc (q), d) /= 0) &
+ stop 5
+
+ if(c_associated (omp_get_mapped_ptr (c_loc (q), d))) &
+ stop 6
+ end if
+
+ call omp_target_free (p, d)
+end program main
diff --git a/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-2.f90 b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-2.f90
new file mode 100644
index 00000000000..66a0b88f612
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-2.f90
@@ -0,0 +1,175 @@
+program main
+ use omp_lib
+ use iso_c_binding
+ implicit none (external, type)
+ integer :: d, id, i, j
+ integer, target :: a, b(1:2), c(1:2), e(0:127)
+ type(c_ptr) :: p1, p2, p3, q, devptrs(0:63)
+
+ a = 42;
+ q = c_loc (a);
+ e = [(i, i = 0, 127)]
+
+ d = omp_get_default_device ()
+ id = omp_get_initial_device ()
+
+ if (d < 0 .or. d >= omp_get_num_devices ()) &
+ d = id
+
+ if (d /= id) then
+ !$omp target data map(alloc: a, b, c(2), e(32:95)) device(d)
+ !$omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c(2), e(32:95)) device(d)
+ p1 = c_loc (a);
+ p2 = c_loc (b);
+ p3 = c_loc (c(2))
+ devptrs = [(c_loc (e(i)), i = 32, 95)]
+ !$omp end target
+
+ if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), p1) &
+ .or. .not. c_associated (omp_get_mapped_ptr (q, d), p1) &
+ .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), p2) &
+ .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d), p2) &
+ .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), p3) &
+ .or. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d))) &
+ stop 0
+
+ do j = 0, 31
+ if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
+ stop 1
+ end do
+ do j = 32, 95
+ if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d), devptrs(j-32))) &
+ stop 2
+ end do
+ do j = 96, 128
+ if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
+ stop 3
+ end do
+ !$omp end target data
+
+ if (c_associated (omp_get_mapped_ptr (c_loc (a), d)) &
+ .or. c_associated (omp_get_mapped_ptr (q, d)) &
+ .or. c_associated (omp_get_mapped_ptr (c_loc (b), d)) &
+ .or. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d)) &
+ .or. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d)) &
+ .or. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d))) &
+ stop 4
+ do j = 0, 127
+ if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
+ stop 5
+ end do
+
+ !$omp target enter data map (alloc: a, b, c(2), e(32:95)) device (d)
+ !$omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c(2), e(32:95)) device(d)
+ p1 = c_loc (a);
+ p2 = c_loc (b);
+ p3 = c_loc (c(2))
+ devptrs = [(c_loc (e(i)), i = 32, 95)]
+ !$omp end target
+
+ if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), p1) &
+ .or. .not. c_associated (omp_get_mapped_ptr (q, d), p1) &
+ .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), p2) &
+ .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), p3) &
+ .or. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d))) &
+ stop 6
+
+ do j = 0, 31
+ if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
+ stop 7
+ end do
+ do j = 32, 95
+ if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d), devptrs(j-32))) &
+ stop 8
+ end do
+ do j = 96, 128
+ if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
+ stop 9
+ end do
+ !$omp target exit data map (delete: a, b, c(2), e(32:95)) device (d)
+
+ if (c_associated (omp_get_mapped_ptr (c_loc (a), d)) &
+ .or. c_associated (omp_get_mapped_ptr (q, d)) &
+ .or. c_associated (omp_get_mapped_ptr (c_loc (b), d)) &
+ .or. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d)) &
+ .or. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d)) &
+ .or. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d))) &
+ stop 10
+ do j = 0, 127
+ if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
+ stop 11
+ end do
+
+ else ! d == id
+
+ !$omp target data map(alloc: a, b, c(2), e(32:95)) device(d)
+ !$omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c(2), e(32:95)) device(d)
+ p1 = c_loc (a);
+ p2 = c_loc (b);
+ p3 = c_loc (c(2))
+ devptrs = [(c_loc (e(i)), i = 32, 95)]
+ !$omp end target
+
+ if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), c_loc (a)) &
+ .or. .not. c_associated (omp_get_mapped_ptr (q, d), q) &
+ .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), c_loc (b)) &
+ .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d), c_loc (b(1))) &
+ .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), c_loc (c(2))) &
+ .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d), c_loc (c(1)))) &
+ stop 12
+
+ do j = 0, 127
+ if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d), c_loc (e(j)))) &
+ stop 13
+ end do
+ !$omp end target data
+
+ if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), c_loc (a)) &
+ .or. .not. c_associated (omp_get_mapped_ptr (q, d), q) &
+ .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), c_loc (b)) &
+ .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d), c_loc (b(1))) &
+ .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), c_loc (c(2))) &
+ .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d), c_loc (c(1)))) &
+ stop 14
+ do j = 0, 127
+ if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
+ stop 15
+ end do
+
+ !$omp target enter data map (alloc: a, b, c(2), e(32:95)) device (d)
+ !$omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c(2), e(32:95)) device(d)
+ p1 = c_loc (a);
+ p2 = c_loc (b);
+ p3 = c_loc (c(2))
+ devptrs = [(c_loc (e(i)), i = 32, 95)]
+ !$omp end target
+
+ if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), c_loc (a)) &
+ .or. .not. c_associated (omp_get_mapped_ptr (q, d), q) &
+ .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), c_loc (b)) &
+ .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d), c_loc (b(1))) &
+ .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), c_loc (c(2))) &
+ .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d), c_loc (c(1)))) &
+ stop 16
+
+ do j = 0, 127
+ if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d), c_loc (e(j)))) &
+ stop 17
+ end do
+ !$omp target exit data map (delete: a, b, c(2), e(32:95)) device (d)
+
+ if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), c_loc (a)) &
+ .or. .not. c_associated (omp_get_mapped_ptr (q, d), q) &
+ .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), c_loc (b)) &
+ .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d), c_loc (b(1))) &
+ .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), c_loc (c(2))) &
+ .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d), c_loc (c(1)))) &
+ stop 18
+
+ do j = 0, 127
+ if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d), c_loc (e(j)))) &
+ stop 19
+ end do
+ end if
+
+end program main
diff --git a/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-3.f90 b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-3.f90
new file mode 100644
index 00000000000..8e7ccac6a52
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-3.f90
@@ -0,0 +1,48 @@
+program main
+ use omp_lib
+ use iso_c_binding
+ implicit none (external, type)
+ integer :: d, id
+ type(c_ptr) :: p, p1, p2
+ integer, target :: a(1:0), b(1:2)
+
+ d = omp_get_default_device ()
+ id = omp_get_initial_device ()
+
+ if (d < 0 .or. d >= omp_get_num_devices ()) &
+ d = id
+
+ p = omp_target_alloc (c_sizeof (c_int), d)
+ if (.not. c_associated (p)) &
+ stop 0 ! okay
+
+ if (omp_target_associate_ptr (c_loc (a), p, c_sizeof (c_int), &
+ 0_c_size_t, d) == 0) then
+
+ if(.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), p)) &
+ stop 1
+
+ if (omp_target_disassociate_ptr (c_loc (a), d) /= 0) &
+ stop 2
+
+ if(c_associated (omp_get_mapped_ptr (c_loc (a), d))) &
+ stop 3
+
+ !$omp target data map(alloc: a) device(d)
+ !$omp target map(from: p1) map(alloc: a) device(d)
+ p1 = c_loc (a);
+ !$omp end target
+ if (c_associated (omp_get_mapped_ptr (c_loc (a), d))) &
+ stop 4
+ !$omp end target data
+
+ !$omp target data map(alloc: b(1:0)) device(d)
+ !$omp target map(from: p2) map(alloc: b(1:0)) device(d)
+ p2 = c_loc (b(1));
+ !$omp end target
+ if (c_associated (omp_get_mapped_ptr (c_loc (b(1)), d))) &
+ stop 5
+ !$omp end target data
+ end if
+ call omp_target_free (p, d)
+end program main
diff --git a/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-4.f90 b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-4.f90
new file mode 100644
index 00000000000..4300a5561ac
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-4.f90
@@ -0,0 +1,84 @@
+program main
+ use omp_lib
+ use iso_c_binding
+ implicit none (external, type)
+ integer :: d, id
+ type(c_ptr) :: p1, p2
+
+ type t
+ integer :: m1, m2
+ end type t
+ type(t), target :: s
+
+ d = omp_get_default_device ()
+ id = omp_get_initial_device ()
+
+ if (d < 0 .or. d >= omp_get_num_devices ()) &
+ d = id
+
+ if (d /= id) then
+ !$omp target data map(alloc: s, s%m2) device(d)
+ !$omp target map(from: p1, p2) map(alloc: s, s%m2) device(d)
+ p1 = c_loc (s);
+ p2 = c_loc (s%m2);
+ !$omp end target
+
+ if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), p1) &
+ .or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), p2)) &
+ stop 0
+ !$omp end target data
+
+ if (c_associated (omp_get_mapped_ptr (c_loc (s), d)) &
+ .or. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d))) &
+ stop 1
+
+ !$omp target enter data map (alloc: s, s%m2) device (d)
+ !$omp target map(from: p1, p2) map(alloc: s, s%m2) device(d)
+ p1 = c_loc (s);
+ p2 = c_loc (s%m2);
+ !$omp end target
+
+ if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), p1) &
+ .or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), p2)) &
+ stop 2
+ !$omp target exit data map (delete: s, s%m2) device (d)
+
+ if (c_associated (omp_get_mapped_ptr (c_loc (s), d)) &
+ .or. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d))) &
+ stop 3
+
+ else ! d == id
+
+ !$omp target data map(alloc: s, s%m2) device(d)
+ !$omp target map(from: p1, p2) map(alloc: s, s%m2) device(d)
+ p1 = c_loc (s);
+ p2 = c_loc (s%m2);
+ !$omp end target
+
+ if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), c_loc (s)) &
+ .or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), c_loc (s%m2))) &
+ stop 4
+ !$omp end target data
+
+ if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), c_loc (s)) &
+ .or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), c_loc (s%m2))) &
+ stop 5
+
+ !$omp target enter data map (alloc: s, s%m2) device (d)
+ !$omp target map(from: p1, p2) map(alloc: s, s%m2) device(d)
+ p1 = c_loc (s);
+ p2 = c_loc (s%m2);
+ !$omp end target
+
+ if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), c_loc (s)) &
+ .or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), c_loc (s%m2))) &
+ stop 6
+
+ !$omp target exit data map (delete: s, s%m2) device (d)
+
+ if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), c_loc (s)) &
+ .or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), c_loc (s%m2))) &
+ stop 7
+ end if
+
+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_get_mapped_ptr 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).