1 #include <omp.h>
2 #include <stdlib.h>
3
4 int a = 1, b = 2, c = 3, d = 4;
5 int e[2] = { 5, 6 }, f[2] = { 7, 8 }, g[2] = { 9, 10 }, h[2] = { 11, 12 };
6
7 __attribute__((noinline, noclone)) void
8 use (int *k, int *l, int *m, int *n, int *o, int *p, int *q, int *r)
9 {
10 asm volatile ("" : : "r" (k) : "memory");
11 asm volatile ("" : : "r" (l) : "memory");
12 asm volatile ("" : : "r" (m) : "memory");
13 asm volatile ("" : : "r" (n) : "memory");
14 asm volatile ("" : : "r" (o) : "memory");
15 asm volatile ("" : : "r" (p) : "memory");
16 asm volatile ("" : : "r" (q) : "memory");
17 asm volatile ("" : : "r" (r) : "memory");
18 }
19
20 #pragma omp declare target to (use)
21
22 int
23 main ()
24 {
25 int err = 0, r = -1, t[4];
26 long s[4] = { -1, -2, -3, -4 };
27 int j = 13, k = 14, l[2] = { 15, 16 }, m[2] = { 17, 18 };
28 #pragma omp target private (a, b, e, f) firstprivate (c, d, g, h) map(from: r, s, t) \
29 map(tofrom: err, j, l) map(to: k, m)
30 #pragma omp teams num_teams (4) thread_limit (8) private (b, f) firstprivate (d, h, k, m)
31 {
32 int u1 = k, u2[2] = { m[0], m[1] };
33 int u3[64];
34 int i;
35 for (i = 0; i < 64; i++)
36 u3[i] = k + i;
37 #pragma omp parallel num_threads (1)
38 {
39 int v1, v2, v3;
40 #pragma omp atomic read
41 v1 = c;
42 #pragma omp atomic read
43 v2 = g[0];
44 #pragma omp atomic read
45 v3 = g[1];
46 if ((v1 < 3 || v1 > 6)
47 || d != 4
48 || (v2 < 9 || v2 > 15 || (v2 & 1) == 0)
49 || (v3 < 10 || v3 > 19 || ((v3 - 10) % 3) != 0)
50 || h[0] != 11 || h[1] != 12 || k != 14 || m[0] != 17 || m[1] != 18)
51 #pragma omp atomic write
52 err = 1;
53 b = omp_get_team_num ();
54 if (b >= 4)
55 #pragma omp atomic write
56 err = 1;
57 if (b == 0)
58 {
59 a = omp_get_num_teams ();
60 e[0] = 2 * a;
61 e[1] = 3 * a;
62 }
63 f[0] = 2 * b;
64 f[1] = 3 * b;
65 #pragma omp atomic update
66 c++;
67 #pragma omp atomic update
68 g[0] += 2;
69 #pragma omp atomic update
70 g[1] += 3;
71 d++;
72 h[0] += 2;
73 h[1] += 3;
74 k += b;
75 m[0] += 2 * b;
76 m[1] += 3 * b;
77 }
78 use (&a, &b, &c, &d, e, f, g, h);
79 #pragma omp parallel firstprivate (u1, u2) private (i)
80 {
81 int w = omp_get_thread_num ();
82 int x = 19;
83 int y[2] = { 20, 21 };
84 int v = 24;
85 int ll[64];
86 if (u1 != 14 || u2[0] != 17 || u2[1] != 18)
87 #pragma omp atomic write
88 err = 1;
89 u1 += w;
90 u2[0] += 2 * w;
91 u2[1] += 3 * w;
92 use (&u1, u2, &t[b], l, &k, m, &j, h);
93 #pragma omp master
94 t[b] = omp_get_num_threads ();
95 #pragma omp atomic update
96 j++;
97 #pragma omp atomic update
98 l[0] += 2;
99 #pragma omp atomic update
100 l[1] += 3;
101 #pragma omp atomic update
102 k += 4;
103 #pragma omp atomic update
104 m[0] += 5;
105 #pragma omp atomic update
106 m[1] += 6;
107 x += w;
108 y[0] += 2 * w;
109 y[1] += 3 * w;
110 #pragma omp simd safelen(32) private (v)
111 for (i = 0; i < 64; i++)
112 {
113 v = 3 * i;
114 ll[i] = u1 + v * u2[0] + u2[1] + x + y[0] + y[1] + v + h[0] + u3[i];
115 }
116 #pragma omp barrier
117 use (&u1, u2, &t[b], l, &k, m, &x, y);
118 if (w < 0 || w > 8 || w != omp_get_thread_num () || u1 != 14 + w
119 || u2[0] != 17 + 2 * w || u2[1] != 18 + 3 * w
120 || x != 19 + w || y[0] != 20 + 2 * w || y[1] != 21 + 3 * w
121 || v != 24)
122 #pragma omp atomic write
123 err = 1;
124 for (i = 0; i < 64; i++)
125 if (ll[i] != u1 + 3 * i * u2[0] + u2[1] + x + y[0] + y[1] + 3 * i + 13 + 14 + i)
126 #pragma omp atomic write
127 err = 1;
128 }
129 #pragma omp parallel num_threads (1)
130 {
131 if (b == 0)
132 {
133 r = a;
134 if (a != omp_get_num_teams ()
135 || e[0] != 2 * a
136 || e[1] != 3 * a)
137 #pragma omp atomic write
138 err = 1;
139 }
140 int v1, v2, v3;
141 #pragma omp atomic read
142 v1 = c;
143 #pragma omp atomic read
144 v2 = g[0];
145 #pragma omp atomic read
146 v3 = g[1];
147 s[b] = v1 * 65536L + v2 * 256L + v3;
148 if (d != 5 || h[0] != 13 || h[1] != 15
149 || k != 14 + b + 4 * t[b]
150 || m[0] != 17 + 2 * b + 5 * t[b]
151 || m[1] != 18 + 3 * b + 6 * t[b]
152 || b != omp_get_team_num ()
153 || f[0] != 2 * b || f[1] != 3 * b)
154 #pragma omp atomic write
155 err = 1;
156 }
157 }
158 if (err != 0) abort ();
159 if (r < 1 || r > 4) abort ();
160 if (a != 1 || b != 2 || c != 3 || d != 4) abort ();
161 if (e[0] != 5 || e[1] != 6 || f[0] != 7 || f[1] != 8) abort ();
162 if (g[0] != 9 || g[1] != 10 || h[0] != 11 || h[1] != 12) abort ();
163 int i, cnt = 0;
164 for (i = 0; i < r; i++)
165 if ((s[i] >> 16) < 3 + 1 || (s[i] >> 16) > 3 + 4
166 || ((s[i] >> 8) & 0xff) < 9 + 2 * 1 || ((s[i] >> 8) & 0xff) > 9 + 2 * 4
167 || (s[i] & 0xff) < 10 + 3 * 1 || (s[i] & 0xff) > 10 + 3 * 4
168 || t[i] < 1 || t[i] > 8)
169 abort ();
170 else
171 cnt += t[i];
172 if (j != 13 + cnt || l[0] != 15 + 2 * cnt || l[1] != 16 + 3 * cnt) abort ();
173 return 0;
174 }