1 /* { dg-do run } */
2
3 #include <omp.h>
4 #include <stdlib.h>
5
6 int v = 6;
7
8 void
9 bar (long *x, long *y)
10 {
11 *x += 2;
12 *y += 3;
13 }
14
15 int
16 baz (void)
17 {
18 return 5;
19 }
20
21 #pragma omp declare target to (bar, baz, v)
22
23 __attribute__((noinline, noclone)) void
24 foo (int a, int b, long c, long d)
25 {
26 int err;
27 if (omp_get_num_teams () != 1)
28 abort ();
29 /* The OpenMP 4.5 spec says that these expressions are evaluated before
30 target region on combined target teams, so those cases are always
31 fine. */
32 #pragma omp target map(from: err)
33 err = omp_get_num_teams () != 1;
34 if (err)
35 abort ();
36 #pragma omp target map(from: err)
37 #pragma omp teams
38 {
39 err = omp_get_num_teams () < 1;
40 #pragma omp parallel if(0)
41 err |= omp_get_thread_limit () < 1;
42 }
43 if (err)
44 abort ();
45 #pragma omp target teams map(from: err)
46 {
47 err = omp_get_num_teams () < 1;
48 #pragma omp parallel if(0)
49 err |= omp_get_thread_limit () < 1;
50 }
51 if (err)
52 abort ();
53 #pragma omp target map(from: err)
54 #pragma omp teams num_teams (4)
55 {
56 err = omp_get_num_teams () < 1 || omp_get_num_teams () > 4;
57 #pragma omp parallel if(0)
58 err |= omp_get_thread_limit () < 1;
59 }
60 if (err)
61 abort ();
62 #pragma omp target teams num_teams (4) map(from: err)
63 {
64 err = omp_get_num_teams () < 1 || omp_get_num_teams () > 4;
65 #pragma omp parallel if(0)
66 err |= omp_get_thread_limit () < 1;
67 }
68 if (err)
69 abort ();
70 #pragma omp target map(from: err)
71 #pragma omp teams thread_limit (7)
72 {
73 err = omp_get_num_teams () < 1;
74 #pragma omp parallel if(0)
75 err |= omp_get_thread_limit () < 1 || omp_get_thread_limit () > 7;
76 }
77 if (err)
78 abort ();
79 #pragma omp target teams thread_limit (7) map(from: err)
80 {
81 err = omp_get_num_teams () < 1;
82 #pragma omp parallel if(0)
83 err |= omp_get_thread_limit () < 1 || omp_get_thread_limit () > 7;
84 }
85 if (err)
86 abort ();
87 #pragma omp target map(from: err)
88 #pragma omp teams num_teams (4) thread_limit (8)
89 {
90 {
91 err = omp_get_num_teams () < 1 || omp_get_num_teams () > 4;
92 }
93 #pragma omp parallel if(0)
94 err |= omp_get_thread_limit () < 1 || omp_get_thread_limit () > 8;
95 }
96 if (err)
97 abort ();
98 #pragma omp target teams num_teams (4) thread_limit (8) map(from: err)
99 {
100 err = omp_get_num_teams () < 1 || omp_get_num_teams () > 4;
101 #pragma omp parallel if(0)
102 err |= omp_get_thread_limit () < 1 || omp_get_thread_limit () > 8;
103 }
104 if (err)
105 abort ();
106 #pragma omp target map(from: err)
107 #pragma omp teams num_teams (a) thread_limit (b)
108 {
109 err = omp_get_num_teams () < 1 || omp_get_num_teams () > a;
110 #pragma omp parallel if(0)
111 err |= omp_get_thread_limit () < 1 || omp_get_thread_limit () > b;
112 }
113 if (err)
114 abort ();
115 #pragma omp target teams num_teams (a) thread_limit (b) map(from: err)
116 {
117 err = omp_get_num_teams () < 1 || omp_get_num_teams () > a;
118 #pragma omp parallel if(0)
119 err |= omp_get_thread_limit () < 1 || omp_get_thread_limit () > b;
120 }
121 if (err)
122 abort ();
123 #pragma omp target map(from: err)
124 #pragma omp teams num_teams (c + 1) thread_limit (d - 1)
125 {
126 err = omp_get_num_teams () < 1 || omp_get_num_teams () > c + 1;
127 #pragma omp parallel if(0)
128 err |= omp_get_thread_limit () < 1 || omp_get_thread_limit () > d - 1;
129 }
130 if (err)
131 abort ();
132 #pragma omp target teams num_teams (c + 1) thread_limit (d - 1) map(from: err)
133 {
134 err = omp_get_num_teams () < 1 || omp_get_num_teams () > c + 1;
135 #pragma omp parallel if(0)
136 err |= omp_get_thread_limit () < 1 || omp_get_thread_limit () > d - 1;
137 }
138 if (err)
139 abort ();
140 #pragma omp target map (always, to: c, d) map(from: err)
141 #pragma omp teams num_teams (c + 1) thread_limit (d - 1)
142 {
143 err = omp_get_num_teams () < 1 || omp_get_num_teams () > c + 1;
144 #pragma omp parallel if(0)
145 err |= omp_get_thread_limit () < 1 || omp_get_thread_limit () > d - 1;
146 }
147 if (err)
148 abort ();
149 #pragma omp target data map (to: c, d)
150 {
151 #pragma omp target defaultmap (tofrom: scalar)
152 bar (&c, &d);
153 /* This is one of the cases which can't be generally optimized,
154 the c and d are (or could be) already mapped and whether
155 their device and original values match is unclear. */
156 #pragma omp target map (to: c, d) map(from: err)
157 #pragma omp teams num_teams (c + 1) thread_limit (d - 1)
158 {
159 err = omp_get_num_teams () < 1 || omp_get_num_teams () > c + 1;
160 #pragma omp parallel if(0)
161 err |= omp_get_thread_limit () < 1 || omp_get_thread_limit () > d - 1;
162 }
163 if (err)
164 abort ();
165 }
166 /* This can't be optimized, there are function calls inside of
167 target involved. */
168 #pragma omp target map(from: err)
169 #pragma omp teams num_teams (baz () + 1) thread_limit (baz () - 1)
170 {
171 err = omp_get_num_teams () < 1 || omp_get_num_teams () > baz () + 1;
172 #pragma omp parallel if(0)
173 err |= omp_get_thread_limit () < 1 || omp_get_thread_limit () > baz () - 1;
174 }
175 if (err)
176 abort ();
177 #pragma omp target teams num_teams (baz () + 1) thread_limit (baz () - 1) map(from: err)
178 {
179 err = omp_get_num_teams () < 1 || omp_get_num_teams () > baz () + 1;
180 #pragma omp parallel if(0)
181 err |= omp_get_thread_limit () < 1 || omp_get_thread_limit () > baz () - 1;
182 }
183 if (err)
184 abort ();
185 /* This one can't be optimized, as v might have different value between
186 host and target. */
187 #pragma omp target map(from: err)
188 #pragma omp teams num_teams (v + 1) thread_limit (v - 1)
189 {
190 err = omp_get_num_teams () < 1 || omp_get_num_teams () > v + 1;
191 #pragma omp parallel if(0)
192 err |= omp_get_thread_limit () < 1 || omp_get_thread_limit () > v - 1;
193 }
194 if (err)
195 abort ();
196 }
197
198 int
199 main ()
200 {
201 foo (3, 5, 7, 9);
202 return 0;
203 }