1 #include <omp.h>
2 #include <stdlib.h>
3
4 struct S { char p[64]; int a; int b[2]; long c[4]; int *d; char q[64]; };
5
6 __attribute__((noinline, noclone)) void
7 foo (struct S s)
8 {
9 int d = omp_get_default_device ();
10 int id = omp_get_initial_device ();
11 int sep = 1;
12
13 if (d < 0 || d >= omp_get_num_devices ())
14 d = id;
15
16 int err;
17 #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d, s.d[-2:3]) map(to: sep) map(from: err)
18 {
19 err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13;
20 err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20;
21 s.a = 35; s.b[0] = 36; s.b[1] = 37;
22 s.c[1] = 38; s.c[2] = 39; s.d[-2] = 40; s.d[-1] = 41; s.d[0] = 42;
23 sep = 0;
24 }
25 if (err) abort ();
26 err = s.a != 35 || s.b[0] != 36 || s.b[1] != 37;
27 err |= s.c[1] != 38 || s.c[2] != 39 || s.d[-2] != 40 || s.d[-1] != 41 || s.d[0] != 42;
28 if (err) abort ();
29 s.a = 50; s.b[0] = 49; s.b[1] = 48;
30 s.c[1] = 47; s.c[2] = 46; s.d[-2] = 45; s.d[-1] = 44; s.d[0] = 43;
31 if (sep
32 && (omp_target_is_present (&s.a, d)
33 || omp_target_is_present (s.b, d)
34 || omp_target_is_present (&s.c[1], d)
35 || omp_target_is_present (s.d, d)
36 || omp_target_is_present (&s.d[-2], d)))
37 abort ();
38 #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
39 {
40 if (!omp_target_is_present (&s.a, d)
41 || !omp_target_is_present (s.b, d)
42 || !omp_target_is_present (&s.c[1], d)
43 || !omp_target_is_present (s.d, d)
44 || !omp_target_is_present (&s.d[-2], d))
45 abort ();
46 #pragma omp target update to(s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
47 #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3]) map(from: err)
48 {
49 err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48;
50 err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43;
51 s.a = 17; s.b[0] = 18; s.b[1] = 19;
52 s.c[1] = 20; s.c[2] = 21; s.d[-2] = 22; s.d[-1] = 23; s.d[0] = 24;
53 }
54 #pragma omp target update from(s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
55 }
56 if (sep
57 && (omp_target_is_present (&s.a, d)
58 || omp_target_is_present (s.b, d)
59 || omp_target_is_present (&s.c[1], d)
60 || omp_target_is_present (s.d, d)
61 || omp_target_is_present (&s.d[-2], d)))
62 abort ();
63 if (err) abort ();
64 err = s.a != 17 || s.b[0] != 18 || s.b[1] != 19;
65 err |= s.c[1] != 20 || s.c[2] != 21 || s.d[-2] != 22 || s.d[-1] != 23 || s.d[0] != 24;
66 if (err) abort ();
67 s.a = 33; s.b[0] = 34; s.b[1] = 35;
68 s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40;
69 #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
70 if (!omp_target_is_present (&s.a, d)
71 || !omp_target_is_present (s.b, d)
72 || !omp_target_is_present (&s.c[1], d)
73 || !omp_target_is_present (s.d, d)
74 || !omp_target_is_present (&s.d[-2], d))
75 abort ();
76 #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
77 #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3]) map(from: err)
78 {
79 err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35;
80 err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40;
81 s.a = 49; s.b[0] = 48; s.b[1] = 47;
82 s.c[1] = 46; s.c[2] = 45; s.d[-2] = 44; s.d[-1] = 43; s.d[0] = 42;
83 }
84 #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
85 if (!omp_target_is_present (&s.a, d)
86 || !omp_target_is_present (s.b, d)
87 || !omp_target_is_present (&s.c[1], d)
88 || !omp_target_is_present (s.d, d)
89 || !omp_target_is_present (&s.d[-2], d))
90 abort ();
91 #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
92 if (sep
93 && (omp_target_is_present (&s.a, d)
94 || omp_target_is_present (s.b, d)
95 || omp_target_is_present (&s.c[1], d)
96 || omp_target_is_present (s.d, d)
97 || omp_target_is_present (&s.d[-2], d)))
98 abort ();
99 if (err) abort ();
100 err = s.a != 49 || s.b[0] != 48 || s.b[1] != 47;
101 err |= s.c[1] != 46 || s.c[2] != 45 || s.d[-2] != 44 || s.d[-1] != 43 || s.d[0] != 42;
102 if (err) abort ();
103 }
104
105 int
106 main ()
107 {
108 int d[3] = { 18, 19, 20 };
109 struct S s = { {}, 11, { 12, 13 }, { 14, 15, 16, 17 }, d + 2, {} };
110 foo (s);
111 return 0;
112 }