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; unsigned char &e; char (&f)[2]; short (&g)[4]; int *&h; char q[64]; };
5 
6 __attribute__((noinline, noclone)) void
foo(S s)7 foo (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[-2:3], s.e, s.f, s.g[1:2], s.h[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     err |= s.e != 21 || s.f[0] != 22 || s.f[1] != 23 || s.g[1] != 25 || s.g[2] != 26;
22     err |= s.h[2] != 31 || s.h[3] != 32 || s.h[4] != 33;
23     s.a = 35; s.b[0] = 36; s.b[1] = 37;
24     s.c[1] = 38; s.c[2] = 39; s.d[-2] = 40; s.d[-1] = 41; s.d[0] = 42;
25     s.e = 43; s.f[0] = 44; s.f[1] = 45; s.g[1] = 46; s.g[2] = 47;
26     s.h[2] = 48; s.h[3] = 49; s.h[4] = 50;
27     sep = 0;
28   }
29   if (err) abort ();
30   err = s.a != 35 || s.b[0] != 36 || s.b[1] != 37;
31   err |= s.c[1] != 38 || s.c[2] != 39 || s.d[-2] != 40 || s.d[-1] != 41 || s.d[0] != 42;
32   err |= s.e != 43 || s.f[0] != 44 || s.f[1] != 45 || s.g[1] != 46 || s.g[2] != 47;
33   err |= s.h[2] != 48 || s.h[3] != 49 || s.h[4] != 50;
34   if (err) abort ();
35   s.a = 50; s.b[0] = 49; s.b[1] = 48;
36   s.c[1] = 47; s.c[2] = 46; s.d[-2] = 45; s.d[-1] = 44; s.d[0] = 43;
37   s.e = 42; s.f[0] = 41; s.f[1] = 40; s.g[1] = 39; s.g[2] = 38;
38   s.h[2] = 37; s.h[3] = 36; s.h[4] = 35;
39   if (sep
40       && (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 	  || omp_target_is_present (&s.e, d)
46 	  || omp_target_is_present (s.f, d)
47 	  || omp_target_is_present (&s.g[1], d)
48 	  || omp_target_is_present (&s.h, d)
49 	  || omp_target_is_present (&s.h[2], d)))
50     abort ();
51   #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
52   {
53     if (!omp_target_is_present (&s.a, d)
54 	|| !omp_target_is_present (s.b, d)
55 	|| !omp_target_is_present (&s.c[1], d)
56 	|| !omp_target_is_present (s.d, d)
57 	|| !omp_target_is_present (&s.d[-2], d)
58 	|| !omp_target_is_present (&s.e, d)
59 	|| !omp_target_is_present (s.f, d)
60 	|| !omp_target_is_present (&s.g[1], d)
61 	|| !omp_target_is_present (&s.h, d)
62 	|| !omp_target_is_present (&s.h[2], d))
63       abort ();
64     #pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
65     #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
66     {
67       err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48;
68       err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43;
69       err |= s.e != 42 || s.f[0] != 41 || s.f[1] != 40 || s.g[1] != 39 || s.g[2] != 38;
70       err |= s.h[2] != 37 || s.h[3] != 36 || s.h[4] != 35;
71       s.a = 17; s.b[0] = 18; s.b[1] = 19;
72       s.c[1] = 20; s.c[2] = 21; s.d[-2] = 22; s.d[-1] = 23; s.d[0] = 24;
73       s.e = 25; s.f[0] = 26; s.f[1] = 27; s.g[1] = 28; s.g[2] = 29;
74       s.h[2] = 30; s.h[3] = 31; s.h[4] = 32;
75     }
76     #pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
77   }
78   if (sep
79       && (omp_target_is_present (&s.a, d)
80 	  || omp_target_is_present (s.b, d)
81 	  || omp_target_is_present (&s.c[1], d)
82 	  || omp_target_is_present (s.d, d)
83 	  || omp_target_is_present (&s.d[-2], d)
84 	  || omp_target_is_present (&s.e, d)
85 	  || omp_target_is_present (s.f, d)
86 	  || omp_target_is_present (&s.g[1], d)
87 	  || omp_target_is_present (&s.h, d)
88 	  || omp_target_is_present (&s.h[2], d)))
89     abort ();
90   if (err) abort ();
91   err = s.a != 17 || s.b[0] != 18 || s.b[1] != 19;
92   err |= s.c[1] != 20 || s.c[2] != 21 || s.d[-2] != 22 || s.d[-1] != 23 || s.d[0] != 24;
93   err |= s.e != 25 || s.f[0] != 26 || s.f[1] != 27 || s.g[1] != 28 || s.g[2] != 29;
94   err |= s.h[2] != 30 || s.h[3] != 31 || s.h[4] != 32;
95   if (err) abort ();
96   s.a = 33; s.b[0] = 34; s.b[1] = 35;
97   s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40;
98   s.e = 41; s.f[0] = 42; s.f[1] = 43; s.g[1] = 44; s.g[2] = 45;
99   s.h[2] = 46; s.h[3] = 47; s.h[4] = 48;
100   #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
101   if (!omp_target_is_present (&s.a, d)
102       || !omp_target_is_present (s.b, d)
103       || !omp_target_is_present (&s.c[1], d)
104       || !omp_target_is_present (s.d, d)
105       || !omp_target_is_present (&s.d[-2], d)
106       || !omp_target_is_present (&s.e, d)
107       || !omp_target_is_present (s.f, d)
108       || !omp_target_is_present (&s.g[1], d)
109       || !omp_target_is_present (&s.h, d)
110       || !omp_target_is_present (&s.h[2], d))
111     abort ();
112   #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
113   #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
114   {
115     err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35;
116     err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40;
117     err |= s.e != 41 || s.f[0] != 42 || s.f[1] != 43 || s.g[1] != 44 || s.g[2] != 45;
118     err |= s.h[2] != 46 || s.h[3] != 47 || s.h[4] != 48;
119     s.a = 49; s.b[0] = 48; s.b[1] = 47;
120     s.c[1] = 46; s.c[2] = 45; s.d[-2] = 44; s.d[-1] = 43; s.d[0] = 42;
121     s.e = 31; s.f[0] = 40; s.f[1] = 39; s.g[1] = 38; s.g[2] = 37;
122     s.h[2] = 36; s.h[3] = 35; s.h[4] = 34;
123   }
124   #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
125   if (!omp_target_is_present (&s.a, d)
126       || !omp_target_is_present (s.b, d)
127       || !omp_target_is_present (&s.c[1], d)
128       || !omp_target_is_present (s.d, d)
129       || !omp_target_is_present (&s.d[-2], d)
130       || !omp_target_is_present (&s.e, d)
131       || !omp_target_is_present (s.f, d)
132       || !omp_target_is_present (&s.g[1], d)
133       || !omp_target_is_present (&s.h, d)
134       || !omp_target_is_present (&s.h[2], d))
135     abort ();
136   #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
137   if (sep
138       && (omp_target_is_present (&s.a, d)
139 	  || omp_target_is_present (s.b, d)
140 	  || omp_target_is_present (&s.c[1], d)
141 	  || omp_target_is_present (s.d, d)
142 	  || omp_target_is_present (&s.d[-2], d)
143 	  || omp_target_is_present (&s.e, d)
144 	  || omp_target_is_present (s.f, d)
145 	  || omp_target_is_present (&s.g[1], d)
146 	  || omp_target_is_present (&s.h, d)
147 	  || omp_target_is_present (&s.h[2], d)))
148     abort ();
149   if (err) abort ();
150   err = s.a != 49 || s.b[0] != 48 || s.b[1] != 47;
151   err |= s.c[1] != 46 || s.c[2] != 45 || s.d[-2] != 44 || s.d[-1] != 43 || s.d[0] != 42;
152   err |= s.e != 31 || s.f[0] != 40 || s.f[1] != 39 || s.g[1] != 38 || s.g[2] != 37;
153   err |= s.h[2] != 36 || s.h[3] != 35 || s.h[4] != 34;
154   if (err) abort ();
155 }
156 
157 int
main()158 main ()
159 {
160   int d[3] = { 18, 19, 20 };
161   unsigned char e = 21;
162   char f[2] = { 22, 23 };
163   short g[4] = { 24, 25, 26, 27 };
164   int hb[7] = { 28, 29, 30, 31, 32, 33, 34 };
165   int *h = hb + 1;
166   S s = { {}, 11, { 12, 13 }, { 14, 15, 16, 17 }, d + 2, e, f, g, h, {} };
167   foo (s);
168 }
169