|
111
|
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
|
|
|
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
|
|
|
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 }
|
