summaryrefslogtreecommitdiffstats
path: root/libgomp
diff options
context:
space:
mode:
authorMarcel Vollweiler <marcel@codesourcery.com>2022-05-02 23:56:44 -0700
committerMarcel Vollweiler <marcel@codesourcery.com>2022-05-02 23:56:44 -0700
commit941cdc8b6d29f9fe494fdd244e96a5e5aa08ba32 (patch)
tree2fd28aefc17121aa268953dace732e22305b3ccf /libgomp
parenttestsuite: vect: update unaligned message (diff)
downloadgcc-941cdc8b6d29f9fe494fdd244e96a5e5aa08ba32.tar.gz
gcc-941cdc8b6d29f9fe494fdd244e96a5e5aa08ba32.tar.bz2
gcc-941cdc8b6d29f9fe494fdd244e96a5e5aa08ba32.tar.xz
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.
Diffstat (limited to 'libgomp')
-rw-r--r--libgomp/libgomp.map5
-rw-r--r--libgomp/libgomp.texi2
-rw-r--r--libgomp/omp.h.in1
-rw-r--r--libgomp/omp_lib.f90.in9
-rw-r--r--libgomp/omp_lib.h.in9
-rw-r--r--libgomp/target.c38
-rw-r--r--libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-1.c41
-rw-r--r--libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-2.c106
-rw-r--r--libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-3.c51
-rw-r--r--libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-4.c49
-rw-r--r--libgomp/testsuite/libgomp.fortran/get-mapped-ptr-1.f9043
-rw-r--r--libgomp/testsuite/libgomp.fortran/get-mapped-ptr-2.f90175
-rw-r--r--libgomp/testsuite/libgomp.fortran/get-mapped-ptr-3.f9048
-rw-r--r--libgomp/testsuite/libgomp.fortran/get-mapped-ptr-4.f9084
14 files changed, 660 insertions, 1 deletions
diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map
index 2ac58094169..608a54cee93 100644
--- a/libgomp/libgomp.map
+++ b/libgomp/libgomp.map
@@ -226,6 +226,11 @@ OMP_5.1 {
226 omp_get_teams_thread_limit_; 226 omp_get_teams_thread_limit_;
227} OMP_5.0.2; 227} OMP_5.0.2;
228 228
229OMP_5.1.1 {
230 global:
231 omp_get_mapped_ptr;
232} OMP_5.1;
233
229GOMP_1.0 { 234GOMP_1.0 {
230 global: 235 global:
231 GOMP_atomic_end; 236 GOMP_atomic_end;
diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
index c10d0cbb6a1..38e0337535a 100644
--- a/libgomp/libgomp.texi
+++ b/libgomp/libgomp.texi
@@ -314,7 +314,7 @@ The OpenMP 4.5 specification is fully supported.
314@item @code{omp_target_is_accessible} runtime routine @tab N @tab 314@item @code{omp_target_is_accessible} runtime routine @tab N @tab
315@item @code{omp_target_memcpy_async} and @code{omp_target_memcpy_rect_async} 315@item @code{omp_target_memcpy_async} and @code{omp_target_memcpy_rect_async}
316 runtime routines @tab N @tab 316 runtime routines @tab N @tab
317@item @code{omp_get_mapped_ptr} runtime routine @tab N @tab 317@item @code{omp_get_mapped_ptr} runtime routine @tab Y @tab
318@item @code{omp_calloc}, @code{omp_realloc}, @code{omp_aligned_alloc} and 318@item @code{omp_calloc}, @code{omp_realloc}, @code{omp_aligned_alloc} and
319 @code{omp_aligned_calloc} runtime routines @tab Y @tab 319 @code{omp_aligned_calloc} runtime routines @tab Y @tab
320@item @code{omp_alloctrait_key_t} enum: @code{omp_atv_serialized} added, 320@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 89c5d65bcd5..18d015295d7 100644
--- a/libgomp/omp.h.in
+++ b/libgomp/omp.h.in
@@ -282,6 +282,7 @@ extern int omp_target_memcpy_rect (void *, const void *, __SIZE_TYPE__, int,
282extern int omp_target_associate_ptr (const void *, const void *, __SIZE_TYPE__, 282extern int omp_target_associate_ptr (const void *, const void *, __SIZE_TYPE__,
283 __SIZE_TYPE__, int) __GOMP_NOTHROW; 283 __SIZE_TYPE__, int) __GOMP_NOTHROW;
284extern int omp_target_disassociate_ptr (const void *, int) __GOMP_NOTHROW; 284extern int omp_target_disassociate_ptr (const void *, int) __GOMP_NOTHROW;
285extern void *omp_get_mapped_ptr (const void *, int) __GOMP_NOTHROW;
285 286
286extern void omp_set_affinity_format (const char *) __GOMP_NOTHROW; 287extern void omp_set_affinity_format (const char *) __GOMP_NOTHROW;
287extern __SIZE_TYPE__ omp_get_affinity_format (char *, __SIZE_TYPE__) 288extern __SIZE_TYPE__ omp_get_affinity_format (char *, __SIZE_TYPE__)
diff --git a/libgomp/omp_lib.f90.in b/libgomp/omp_lib.f90.in
index daf40dc8e6d..506f15cdaae 100644
--- a/libgomp/omp_lib.f90.in
+++ b/libgomp/omp_lib.f90.in
@@ -835,6 +835,15 @@
835 end function omp_target_disassociate_ptr 835 end function omp_target_disassociate_ptr
836 end interface 836 end interface
837 837
838 interface
839 function omp_get_mapped_ptr (ptr, device_num) bind(c)
840 use, intrinsic :: iso_c_binding, only : c_ptr, c_int
841 type(c_ptr) :: omp_get_mapped_ptr
842 type(c_ptr), value :: ptr
843 integer(c_int), value :: device_num
844 end function omp_get_mapped_ptr
845 end interface
846
838#if _OPENMP >= 201811 847#if _OPENMP >= 201811
839!GCC$ ATTRIBUTES DEPRECATED :: omp_get_nested, omp_set_nested 848!GCC$ ATTRIBUTES DEPRECATED :: omp_get_nested, omp_set_nested
840#endif 849#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 @@
416 integer(c_int), value :: device_num 416 integer(c_int), value :: device_num
417 end function omp_target_disassociate_ptr 417 end function omp_target_disassociate_ptr
418 end interface 418 end interface
419
420 interface
421 function omp_get_mapped_ptr (ptr, device_num) bind(c)
422 use, intrinsic :: iso_c_binding, only : c_ptr, c_int
423 type(c_ptr) :: omp_get_mapped_ptr
424 type(c_ptr), value :: ptr
425 integer(c_int), value :: device_num
426 end function omp_get_mapped_ptr
427 end interface
diff --git a/libgomp/target.c b/libgomp/target.c
index 9017458885e..86930ea2d5d 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -3665,6 +3665,44 @@ omp_target_disassociate_ptr (const void *ptr, int device_num)
3665 return ret; 3665 return ret;
3666} 3666}
3667 3667
3668void *
3669omp_get_mapped_ptr (const void *ptr, int device_num)
3670{
3671 if (device_num < 0 || device_num > gomp_get_num_devices ())
3672 return NULL;
3673
3674 if (device_num == omp_get_initial_device ())
3675 return (void *) ptr;
3676
3677 struct gomp_device_descr *devicep = resolve_device (device_num);
3678 if (devicep == NULL)
3679 return NULL;
3680
3681 if (!(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
3682 || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
3683 return (void *) ptr;
3684
3685 gomp_mutex_lock (&devicep->lock);
3686
3687 struct splay_tree_s *mem_map = &devicep->mem_map;
3688 struct splay_tree_key_s cur_node;
3689 void *ret = NULL;
3690
3691 cur_node.host_start = (uintptr_t) ptr;
3692 cur_node.host_end = cur_node.host_start;
3693 splay_tree_key n = gomp_map_0len_lookup (mem_map, &cur_node);
3694
3695 if (n)
3696 {
3697 uintptr_t offset = cur_node.host_start - n->host_start;
3698 ret = (void *) (n->tgt->tgt_start + n->tgt_offset + offset);
3699 }
3700
3701 gomp_mutex_unlock (&devicep->lock);
3702
3703 return ret;
3704}
3705
3668int 3706int
3669omp_pause_resource (omp_pause_resource_t kind, int device_num) 3707omp_pause_resource (omp_pause_resource_t kind, int device_num)
3670{ 3708{
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 @@
1#include <omp.h>
2#include <stdlib.h>
3
4int
5main ()
6{
7 int d = omp_get_default_device ();
8 int id = omp_get_initial_device ();
9 void *p , *q;
10
11 if (d < 0 || d >= omp_get_num_devices ())
12 d = id;
13
14 p = omp_target_alloc (sizeof (int), d);
15 if (p == NULL)
16 return 0;
17
18 if (omp_target_associate_ptr (q, p, sizeof (int), 0, d) != 0)
19 return 0;
20
21 if (omp_get_mapped_ptr (q, -1) != NULL)
22 abort ();
23
24 if (omp_get_mapped_ptr (q, omp_get_num_devices () + 1) != NULL)
25 abort ();
26
27 if (omp_get_mapped_ptr (q, id) != q)
28 abort ();
29
30 if (omp_get_mapped_ptr (q, d) != p)
31 abort ();
32
33 if (omp_target_disassociate_ptr (q, d) != 0)
34 abort ();
35
36 if (omp_get_mapped_ptr (q, d) != NULL)
37 abort ();
38
39 omp_target_free (p, d);
40 return 0;
41}
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 @@
1#include <omp.h>
2#include <stdlib.h>
3#include <stdint.h>
4
5int
6main ()
7{
8 int d = omp_get_default_device ();
9 int id = omp_get_initial_device ();
10 int a = 42;
11 int b[] = { 24, 42 };
12 int c[] = { 47, 11 };
13 int e[128];
14 int *q = &a;
15 void *p1 = NULL, *p2 = NULL, *p3 = NULL;
16 void *devptrs[128];
17
18 if (d < 0 || d >= omp_get_num_devices ())
19 d = id;
20
21 for (int i = 0; i < 128; i++)
22 e[i] = i;
23
24 #pragma omp target data map(alloc: a, b, c[1], e[32:64]) device(d)
25 {
26 #pragma omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c[1], e[32:64]) device(d)
27 {
28 p1 = &a;
29 p2 = &b;
30 p3 = &c[1];
31 for (int i = 32; i < 96; i++)
32 devptrs[i] = &e[i];
33 }
34
35 if (omp_get_mapped_ptr (&a, d) != (d == id ? &a : p1)
36 || omp_get_mapped_ptr (q, d) != (d == id ? q : p1)
37 || omp_get_mapped_ptr (b, d) != (d == id ? b : p2)
38 || omp_get_mapped_ptr (&b[0], d) != (d == id ? &b[0] : p2)
39 || omp_get_mapped_ptr (&c[1], d) != (d == id ? &c[1] : p3)
40 || omp_get_mapped_ptr (&c[0], d) != (d == id ? &c[0] : NULL))
41 abort ();
42
43 for (int i = 0; i < 32; i++)
44 if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : NULL))
45 abort ();
46 for (int i = 32; i < 96; i++)
47 if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : devptrs[i]))
48 abort ();
49 for (int i = 96; i < 128; i++)
50 if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : NULL))
51 abort ();
52 }
53
54 if (omp_get_mapped_ptr (&a, d) != (d == id ? &a : NULL)
55 || omp_get_mapped_ptr (q, d) != (d == id ? q : NULL)
56 || omp_get_mapped_ptr (b, d) != (d == id ? b : NULL)
57 || omp_get_mapped_ptr (&b[0], d) != (d == id ? &b[0] : NULL)
58 || omp_get_mapped_ptr (&c[1], d) != (d == id ? &c[1] : NULL)
59 || omp_get_mapped_ptr (&c[0], d) != (d == id ? &c[0] : NULL))
60 abort ();
61 for (int i = 0; i < 128; i++)
62 if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : NULL))
63 abort ();
64
65 #pragma omp target enter data map (alloc: a, b, c[1], e[32:64]) device (d)
66 #pragma omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c[1], e[32:64]) device(d)
67 {
68 p1 = &a;
69 p2 = &b;
70 p3 = &c[1];
71 for (int i = 32; i < 96; i++)
72 devptrs[i] = &e[i];
73 }
74
75 if (omp_get_mapped_ptr (&a, d) != (d == id ? &a : p1)
76 || omp_get_mapped_ptr (q, d) != (d == id ? q : p1)
77 || omp_get_mapped_ptr (b, d) != (d == id ? b : p2)
78 || omp_get_mapped_ptr (&b[0], d) != (d == id ? &b[0] : p2)
79 || omp_get_mapped_ptr (&c[1], d) != (d == id ? &c[1] : p3)
80 || omp_get_mapped_ptr (&c[0], d) != (d == id ? &c[0] : NULL))
81 abort ();
82 for (int i = 0; i < 32; i++)
83 if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : NULL))
84 abort ();
85 for (int i = 32; i < 96; i++)
86 if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : devptrs[i]))
87 abort ();
88 for (int i = 96; i < 128; i++)
89 if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : NULL))
90 abort ();
91
92 #pragma omp target exit data map (delete: a, b, c[1], e[32:64]) device (d)
93
94 if (omp_get_mapped_ptr (&a, d) != (d == id ? &a : NULL)
95 || omp_get_mapped_ptr (q, d) != (d == id ? q : NULL)
96 || omp_get_mapped_ptr (b, d) != (d == id ? b : NULL)
97 || omp_get_mapped_ptr (&b[0], d) != (d == id ? &b[0] : NULL)
98 || omp_get_mapped_ptr (&c[1], d) != (d == id ? &c[1] : NULL)
99 || omp_get_mapped_ptr (&c[0], d) != (d == id ? &c[0] : NULL))
100 abort ();
101 for (int i = 0; i < 128; i++)
102 if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : NULL))
103 abort ();
104
105 return 0;
106}
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 @@
1#include <omp.h>
2#include <stdlib.h>
3
4int
5main ()
6{
7 int d = omp_get_default_device ();
8 int id = omp_get_initial_device ();
9 int a[0];
10 int b[] = { 24, 42 };
11 void *p1 = NULL, *p2 = NULL;
12
13 if (d < 0 || d >= omp_get_num_devices ())
14 d = id;
15
16 void *p = omp_target_alloc (sizeof (int), d);
17 if (p == NULL)
18 return 0;
19
20 if (omp_target_associate_ptr (a, p, sizeof (int), 0, d) != 0)
21 return 0;
22
23 if (omp_get_mapped_ptr (a, d) != (d == id ? a : p))
24 abort ();
25
26 if (omp_target_disassociate_ptr (a, d) != 0)
27 abort ();
28
29 if (omp_get_mapped_ptr (a, d) != (d == id ? a : NULL))
30 abort ();
31
32 #pragma omp target data map(alloc: a, b[1:0]) device(d)
33 {
34 #pragma omp target map(from: p1, p2) map(alloc: a, b[1:0]) device(d)
35 {
36 p1 = &a;
37 p2 = &b[1];
38 }
39
40 /* This is probably expected to be p1/p2 instead of NULL. Zero-length arrays
41 as list items of the map clause are currently not inserted into the mem
42 map ?! However by returning NULL, omp_get_mapped_ptr is consistent with
43 omp_target_is_present. */
44 if (omp_get_mapped_ptr (a, d) != NULL
45 || omp_get_mapped_ptr (&b[1], d) != NULL)
46 abort ();
47 }
48
49 omp_target_free (p, d);
50 return 0;
51}
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 @@
1#include <omp.h>
2#include <stdlib.h>
3
4int
5main ()
6{
7 int d = omp_get_default_device ();
8 int id = omp_get_initial_device ();
9 struct s_t { int m1; char m2; } s;
10 void *p1 = NULL, *p2 = NULL;
11
12 if (d < 0 || d >= omp_get_num_devices ())
13 d = id;
14
15 #pragma omp target data map(alloc: s, s.m2) device(d)
16 {
17 #pragma omp target map(from: p1, p2) map(alloc: s, s.m2) device(d)
18 {
19 p1 = &s;
20 p2 = &s.m2;
21 }
22 if (omp_get_mapped_ptr (&s, d) != (d == id ? &s : p1)
23 || omp_get_mapped_ptr (&s.m2, d) != (d == id ? &s.m2 : p2))
24 abort ();
25 }
26
27 if (omp_get_mapped_ptr (&s, d) != (d == id ? &s : NULL)
28 || omp_get_mapped_ptr (&s.m2, d) != (d == id ? &s.m2 : NULL))
29 abort ();
30
31 #pragma omp target enter data map(alloc: s, s.m2) device (d)
32 #pragma omp target map(from: p1, p2) map(alloc: s, s.m2) device(d)
33 {
34 p1 = &s;
35 p2 = &s.m2;
36 }
37
38 if (omp_get_mapped_ptr (&s, d) != (d == id ? &s : p1)
39 || omp_get_mapped_ptr (&s.m2, d) != (d == id ? &s.m2 : p2))
40 abort ();
41
42 #pragma omp target exit data map (delete: s, s.m2) device (d)
43
44 if (omp_get_mapped_ptr (&s, d) != (d == id ? &s : NULL)
45 || omp_get_mapped_ptr (&s.m2, d) != (d == id ? &s.m2 : NULL))
46 abort ();
47
48 return 0;
49}
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 @@
1program main
2 use omp_lib
3 use iso_c_binding
4 implicit none (external, type)
5 integer :: d, id
6 type(c_ptr) :: p
7 integer, target :: q
8
9 d = omp_get_default_device ()
10 id = omp_get_initial_device ()
11
12 if (d < 0 .or. d >= omp_get_num_devices ()) &
13 d = id
14
15 p = omp_target_alloc (c_sizeof (q), d)
16 if (.not. c_associated (p)) &
17 stop 0 ! okay
18
19 if (omp_target_associate_ptr (c_loc (q), p, c_sizeof (q), &
20 0_c_size_t, d) == 0) then
21
22 if(c_associated (omp_get_mapped_ptr (c_loc (q), -1))) &
23 stop 1
24
25 if(c_associated (omp_get_mapped_ptr (c_loc (q), &
26 omp_get_num_devices () + 1))) &
27 stop 2
28
29 if(.not. c_associated (omp_get_mapped_ptr (c_loc (q), id), c_loc (q))) &
30 stop 3
31
32 if(.not. c_associated (omp_get_mapped_ptr (c_loc (q), d), p)) &
33 stop 4
34
35 if (omp_target_disassociate_ptr (c_loc (q), d) /= 0) &
36 stop 5
37
38 if(c_associated (omp_get_mapped_ptr (c_loc (q), d))) &
39 stop 6
40 end if
41
42 call omp_target_free (p, d)
43end 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 @@
1program main
2 use omp_lib
3 use iso_c_binding
4 implicit none (external, type)
5 integer :: d, id, i, j
6 integer, target :: a, b(1:2), c(1:2), e(0:127)
7 type(c_ptr) :: p1, p2, p3, q, devptrs(0:63)
8
9 a = 42;
10 q = c_loc (a);
11 e = [(i, i = 0, 127)]
12
13 d = omp_get_default_device ()
14 id = omp_get_initial_device ()
15
16 if (d < 0 .or. d >= omp_get_num_devices ()) &
17 d = id
18
19 if (d /= id) then
20 !$omp target data map(alloc: a, b, c(2), e(32:95)) device(d)
21 !$omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c(2), e(32:95)) device(d)
22 p1 = c_loc (a);
23 p2 = c_loc (b);
24 p3 = c_loc (c(2))
25 devptrs = [(c_loc (e(i)), i = 32, 95)]
26 !$omp end target
27
28 if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), p1) &
29 .or. .not. c_associated (omp_get_mapped_ptr (q, d), p1) &
30 .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), p2) &
31 .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d), p2) &
32 .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), p3) &
33 .or. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d))) &
34 stop 0
35
36 do j = 0, 31
37 if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
38 stop 1
39 end do
40 do j = 32, 95
41 if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d), devptrs(j-32))) &
42 stop 2
43 end do
44 do j = 96, 128
45 if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
46 stop 3
47 end do
48 !$omp end target data
49
50 if (c_associated (omp_get_mapped_ptr (c_loc (a), d)) &
51 .or. c_associated (omp_get_mapped_ptr (q, d)) &
52 .or. c_associated (omp_get_mapped_ptr (c_loc (b), d)) &
53 .or. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d)) &
54 .or. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d)) &
55 .or. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d))) &
56 stop 4
57 do j = 0, 127
58 if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
59 stop 5
60 end do
61
62 !$omp target enter data map (alloc: a, b, c(2), e(32:95)) device (d)
63 !$omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c(2), e(32:95)) device(d)
64 p1 = c_loc (a);
65 p2 = c_loc (b);
66 p3 = c_loc (c(2))
67 devptrs = [(c_loc (e(i)), i = 32, 95)]
68 !$omp end target
69
70 if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), p1) &
71 .or. .not. c_associated (omp_get_mapped_ptr (q, d), p1) &
72 .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), p2) &
73 .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), p3) &
74 .or. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d))) &
75 stop 6
76
77 do j = 0, 31
78 if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
79 stop 7
80 end do
81 do j = 32, 95
82 if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d), devptrs(j-32))) &
83 stop 8
84 end do
85 do j = 96, 128
86 if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
87 stop 9
88 end do
89 !$omp target exit data map (delete: a, b, c(2), e(32:95)) device (d)
90
91 if (c_associated (omp_get_mapped_ptr (c_loc (a), d)) &
92 .or. c_associated (omp_get_mapped_ptr (q, d)) &
93 .or. c_associated (omp_get_mapped_ptr (c_loc (b), d)) &
94 .or. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d)) &
95 .or. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d)) &
96 .or. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d))) &
97 stop 10
98 do j = 0, 127
99 if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
100 stop 11
101 end do
102
103 else ! d == id
104
105 !$omp target data map(alloc: a, b, c(2), e(32:95)) device(d)
106 !$omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c(2), e(32:95)) device(d)
107 p1 = c_loc (a);
108 p2 = c_loc (b);
109 p3 = c_loc (c(2))
110 devptrs = [(c_loc (e(i)), i = 32, 95)]
111 !$omp end target
112
113 if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), c_loc (a)) &
114 .or. .not. c_associated (omp_get_mapped_ptr (q, d), q) &
115 .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), c_loc (b)) &
116 .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d), c_loc (b(1))) &
117 .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), c_loc (c(2))) &
118 .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d), c_loc (c(1)))) &
119 stop 12
120
121 do j = 0, 127
122 if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d), c_loc (e(j)))) &
123 stop 13
124 end do
125 !$omp end target data
126
127 if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), c_loc (a)) &
128 .or. .not. c_associated (omp_get_mapped_ptr (q, d), q) &
129 .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), c_loc (b)) &
130 .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d), c_loc (b(1))) &
131 .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), c_loc (c(2))) &
132 .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d), c_loc (c(1)))) &
133 stop 14
134 do j = 0, 127
135 if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) &
136 stop 15
137 end do
138
139 !$omp target enter data map (alloc: a, b, c(2), e(32:95)) device (d)
140 !$omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c(2), e(32:95)) device(d)
141 p1 = c_loc (a);
142 p2 = c_loc (b);
143 p3 = c_loc (c(2))
144 devptrs = [(c_loc (e(i)), i = 32, 95)]
145 !$omp end target
146
147 if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), c_loc (a)) &
148 .or. .not. c_associated (omp_get_mapped_ptr (q, d), q) &
149 .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), c_loc (b)) &
150 .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d), c_loc (b(1))) &
151 .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), c_loc (c(2))) &
152 .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d), c_loc (c(1)))) &
153 stop 16
154
155 do j = 0, 127
156 if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d), c_loc (e(j)))) &
157 stop 17
158 end do
159 !$omp target exit data map (delete: a, b, c(2), e(32:95)) device (d)
160
161 if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), c_loc (a)) &
162 .or. .not. c_associated (omp_get_mapped_ptr (q, d), q) &
163 .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), c_loc (b)) &
164 .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d), c_loc (b(1))) &
165 .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), c_loc (c(2))) &
166 .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d), c_loc (c(1)))) &
167 stop 18
168
169 do j = 0, 127
170 if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d), c_loc (e(j)))) &
171 stop 19
172 end do
173 end if
174
175end 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 @@
1program main
2 use omp_lib
3 use iso_c_binding
4 implicit none (external, type)
5 integer :: d, id
6 type(c_ptr) :: p, p1, p2
7 integer, target :: a(1:0), b(1:2)
8
9 d = omp_get_default_device ()
10 id = omp_get_initial_device ()
11
12 if (d < 0 .or. d >= omp_get_num_devices ()) &
13 d = id
14
15 p = omp_target_alloc (c_sizeof (c_int), d)
16 if (.not. c_associated (p)) &
17 stop 0 ! okay
18
19 if (omp_target_associate_ptr (c_loc (a), p, c_sizeof (c_int), &
20 0_c_size_t, d) == 0) then
21
22 if(.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), p)) &
23 stop 1
24
25 if (omp_target_disassociate_ptr (c_loc (a), d) /= 0) &
26 stop 2
27
28 if(c_associated (omp_get_mapped_ptr (c_loc (a), d))) &
29 stop 3
30
31 !$omp target data map(alloc: a) device(d)
32 !$omp target map(from: p1) map(alloc: a) device(d)
33 p1 = c_loc (a);
34 !$omp end target
35 if (c_associated (omp_get_mapped_ptr (c_loc (a), d))) &
36 stop 4
37 !$omp end target data
38
39 !$omp target data map(alloc: b(1:0)) device(d)
40 !$omp target map(from: p2) map(alloc: b(1:0)) device(d)
41 p2 = c_loc (b(1));
42 !$omp end target
43 if (c_associated (omp_get_mapped_ptr (c_loc (b(1)), d))) &
44 stop 5
45 !$omp end target data
46 end if
47 call omp_target_free (p, d)
48end 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 @@
1program main
2 use omp_lib
3 use iso_c_binding
4 implicit none (external, type)
5 integer :: d, id
6 type(c_ptr) :: p1, p2
7
8 type t
9 integer :: m1, m2
10 end type t
11 type(t), target :: s
12
13 d = omp_get_default_device ()
14 id = omp_get_initial_device ()
15
16 if (d < 0 .or. d >= omp_get_num_devices ()) &
17 d = id
18
19 if (d /= id) then
20 !$omp target data map(alloc: s, s%m2) device(d)
21 !$omp target map(from: p1, p2) map(alloc: s, s%m2) device(d)
22 p1 = c_loc (s);
23 p2 = c_loc (s%m2);
24 !$omp end target
25
26 if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), p1) &
27 .or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), p2)) &
28 stop 0
29 !$omp end target data
30
31 if (c_associated (omp_get_mapped_ptr (c_loc (s), d)) &
32 .or. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d))) &
33 stop 1
34
35 !$omp target enter data map (alloc: s, s%m2) device (d)
36 !$omp target map(from: p1, p2) map(alloc: s, s%m2) device(d)
37 p1 = c_loc (s);
38 p2 = c_loc (s%m2);
39 !$omp end target
40
41 if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), p1) &
42 .or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), p2)) &
43 stop 2
44 !$omp target exit data map (delete: s, s%m2) device (d)
45
46 if (c_associated (omp_get_mapped_ptr (c_loc (s), d)) &
47 .or. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d))) &
48 stop 3
49
50 else ! d == id
51
52 !$omp target data map(alloc: s, s%m2) device(d)
53 !$omp target map(from: p1, p2) map(alloc: s, s%m2) device(d)
54 p1 = c_loc (s);
55 p2 = c_loc (s%m2);
56 !$omp end target
57
58 if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), c_loc (s)) &
59 .or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), c_loc (s%m2))) &
60 stop 4
61 !$omp end target data
62
63 if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), c_loc (s)) &
64 .or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), c_loc (s%m2))) &
65 stop 5
66
67 !$omp target enter data map (alloc: s, s%m2) device (d)
68 !$omp target map(from: p1, p2) map(alloc: s, s%m2) device(d)
69 p1 = c_loc (s);
70 p2 = c_loc (s%m2);
71 !$omp end target
72
73 if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), c_loc (s)) &
74 .or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), c_loc (s%m2))) &
75 stop 6
76
77 !$omp target exit data map (delete: s, s%m2) device (d)
78
79 if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), c_loc (s)) &
80 .or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), c_loc (s%m2))) &
81 stop 7
82 end if
83
84end program main