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