Line data Source code
1 : /* General types and functions that are useful for processing of OpenMP,
2 : OpenACC and similar directives at various stages of compilation.
3 :
4 : Copyright (C) 2005-2026 Free Software Foundation, Inc.
5 :
6 : This file is part of GCC.
7 :
8 : GCC is free software; you can redistribute it and/or modify it under
9 : the terms of the GNU General Public License as published by the Free
10 : Software Foundation; either version 3, or (at your option) any later
11 : version.
12 :
13 : GCC is distributed in the hope that it will be useful, but WITHOUT ANY
14 : WARRANTY; without even the implied warranty of MERCHANTABILITY or
15 : FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
16 : for more details.
17 :
18 : You should have received a copy of the GNU General Public License
19 : along with GCC; see the file COPYING3. If not see
20 : <http://www.gnu.org/licenses/>. */
21 :
22 : #include "config.h"
23 : #include "system.h"
24 : #include "coretypes.h"
25 : #include "backend.h"
26 : #include "target.h"
27 : #include "tree.h"
28 : #include "gimple.h"
29 : #include "ssa.h"
30 : #include "diagnostic-core.h"
31 : #include "fold-const.h"
32 : #include "langhooks.h"
33 : #include "omp-general.h"
34 : #include "stringpool.h"
35 : #include "attribs.h"
36 : #include "gimplify.h"
37 : #include "cgraph.h"
38 : #include "alloc-pool.h"
39 : #include "symbol-summary.h"
40 : #include "tree-pass.h"
41 : #include "omp-device-properties.h"
42 : #include "tree-iterator.h"
43 : #include "data-streamer.h"
44 : #include "streamer-hooks.h"
45 : #include "opts.h"
46 : #include "tree-pretty-print.h"
47 :
48 : enum omp_requires omp_requires_mask;
49 :
50 : /* Find an OMP clause of type KIND within CLAUSES. */
51 : tree
52 1222193 : omp_find_clause (tree clauses, enum omp_clause_code kind)
53 : {
54 4494026 : for (; clauses ; clauses = OMP_CLAUSE_CHAIN (clauses))
55 3528400 : if (OMP_CLAUSE_CODE (clauses) == kind)
56 : return clauses;
57 :
58 : return NULL_TREE;
59 : }
60 :
61 : /* True if OpenMP should regard this DECL as being a scalar which has Fortran's
62 : allocatable or pointer attribute. */
63 : bool
64 15434 : omp_is_allocatable_or_ptr (tree decl)
65 : {
66 15434 : return lang_hooks.decls.omp_is_allocatable_or_ptr (decl);
67 : }
68 :
69 : /* Check whether this DECL belongs to a Fortran optional argument.
70 : With 'for_present_check' set to false, decls which are optional parameters
71 : themselves are returned as tree - or a NULL_TREE otherwise. Those decls are
72 : always pointers. With 'for_present_check' set to true, the decl for checking
73 : whether an argument is present is returned; for arguments with value
74 : attribute this is the hidden argument and of BOOLEAN_TYPE. If the decl is
75 : unrelated to optional arguments, NULL_TREE is returned. */
76 :
77 : tree
78 13288 : omp_check_optional_argument (tree decl, bool for_present_check)
79 : {
80 13288 : return lang_hooks.decls.omp_check_optional_argument (decl, for_present_check);
81 : }
82 :
83 : /* Return true if TYPE is an OpenMP mappable type. */
84 :
85 : bool
86 58765 : omp_mappable_type (tree type)
87 : {
88 : /* Mappable type has to be complete. */
89 58765 : if (type == error_mark_node || !COMPLETE_TYPE_P (type))
90 190 : return false;
91 : return true;
92 : }
93 :
94 : /* True if OpenMP should privatize what this DECL points to rather
95 : than the DECL itself. */
96 :
97 : bool
98 5026635 : omp_privatize_by_reference (tree decl)
99 : {
100 5026635 : return lang_hooks.decls.omp_privatize_by_reference (decl);
101 : }
102 :
103 : /* Adjust *COND_CODE and *N2 so that the former is either LT_EXPR or GT_EXPR,
104 : given that V is the loop index variable and STEP is loop step. */
105 :
106 : void
107 196508 : omp_adjust_for_condition (location_t loc, enum tree_code *cond_code, tree *n2,
108 : tree v, tree step)
109 : {
110 196508 : switch (*cond_code)
111 : {
112 : case LT_EXPR:
113 : case GT_EXPR:
114 : break;
115 :
116 29865 : case NE_EXPR:
117 29865 : gcc_assert (TREE_CODE (step) == INTEGER_CST);
118 29865 : if (TREE_CODE (TREE_TYPE (v)) == INTEGER_TYPE
119 29865 : || TREE_CODE (TREE_TYPE (v)) == BITINT_TYPE)
120 : {
121 25021 : if (integer_onep (step))
122 17737 : *cond_code = LT_EXPR;
123 : else
124 : {
125 7284 : gcc_assert (integer_minus_onep (step));
126 7284 : *cond_code = GT_EXPR;
127 : }
128 : }
129 : else
130 : {
131 4844 : tree unit = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (v)));
132 4844 : gcc_assert (TREE_CODE (unit) == INTEGER_CST);
133 4844 : if (tree_int_cst_equal (unit, step))
134 3358 : *cond_code = LT_EXPR;
135 : else
136 : {
137 1486 : gcc_assert (wi::neg (wi::to_widest (unit))
138 : == wi::to_widest (step));
139 1486 : *cond_code = GT_EXPR;
140 : }
141 : }
142 :
143 : break;
144 :
145 28283 : case LE_EXPR:
146 28283 : if (POINTER_TYPE_P (TREE_TYPE (*n2)))
147 : {
148 123 : tree unit = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (v)));
149 123 : gcc_assert (TREE_CODE (unit) == INTEGER_CST);
150 123 : *n2 = fold_build_pointer_plus_loc (loc, *n2, unit);
151 : }
152 : else
153 28160 : *n2 = fold_build2_loc (loc, PLUS_EXPR, TREE_TYPE (*n2), *n2,
154 28160 : build_int_cst (TREE_TYPE (*n2), 1));
155 28283 : *cond_code = LT_EXPR;
156 28283 : break;
157 1808 : case GE_EXPR:
158 1808 : if (POINTER_TYPE_P (TREE_TYPE (*n2)))
159 : {
160 137 : tree unit = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (v)));
161 137 : gcc_assert (TREE_CODE (unit) == INTEGER_CST);
162 137 : unit = convert_to_ptrofftype_loc (loc, unit);
163 137 : unit = fold_build1_loc (loc, NEGATE_EXPR, TREE_TYPE (unit),
164 : unit);
165 137 : *n2 = fold_build_pointer_plus_loc (loc, *n2, unit);
166 : }
167 : else
168 1671 : *n2 = fold_build2_loc (loc, MINUS_EXPR, TREE_TYPE (*n2), *n2,
169 1671 : build_int_cst (TREE_TYPE (*n2), 1));
170 1808 : *cond_code = GT_EXPR;
171 1808 : break;
172 0 : default:
173 0 : gcc_unreachable ();
174 : }
175 196508 : }
176 :
177 : /* Return the looping step from INCR, extracted from the step of a gimple omp
178 : for statement. */
179 :
180 : tree
181 195344 : omp_get_for_step_from_incr (location_t loc, tree incr)
182 : {
183 195344 : tree step;
184 195344 : switch (TREE_CODE (incr))
185 : {
186 173096 : case PLUS_EXPR:
187 173096 : step = TREE_OPERAND (incr, 1);
188 173096 : break;
189 16144 : case POINTER_PLUS_EXPR:
190 16144 : step = fold_convert (ssizetype, TREE_OPERAND (incr, 1));
191 16144 : break;
192 6104 : case MINUS_EXPR:
193 6104 : step = TREE_OPERAND (incr, 1);
194 6104 : step = fold_build1_loc (loc, NEGATE_EXPR, TREE_TYPE (step), step);
195 6104 : break;
196 0 : default:
197 0 : gcc_unreachable ();
198 : }
199 195344 : return step;
200 : }
201 :
202 : /* Extract the header elements of parallel loop FOR_STMT and store
203 : them into *FD. */
204 :
205 : void
206 129291 : omp_extract_for_data (gomp_for *for_stmt, struct omp_for_data *fd,
207 : struct omp_for_data_loop *loops)
208 : {
209 129291 : tree t, var, *collapse_iter, *collapse_count;
210 129291 : tree count = NULL_TREE, iter_type = long_integer_type_node;
211 129291 : struct omp_for_data_loop *loop;
212 129291 : int i;
213 129291 : struct omp_for_data_loop dummy_loop;
214 129291 : location_t loc = gimple_location (for_stmt);
215 129291 : bool simd = gimple_omp_for_kind (for_stmt) == GF_OMP_FOR_KIND_SIMD;
216 129291 : bool distribute = gimple_omp_for_kind (for_stmt)
217 129291 : == GF_OMP_FOR_KIND_DISTRIBUTE;
218 129291 : bool taskloop = gimple_omp_for_kind (for_stmt)
219 129291 : == GF_OMP_FOR_KIND_TASKLOOP;
220 129291 : bool order_reproducible = false;
221 129291 : tree iterv, countv;
222 :
223 129291 : fd->for_stmt = for_stmt;
224 129291 : fd->pre = NULL;
225 129291 : fd->have_nowait = distribute || simd;
226 129291 : fd->have_ordered = false;
227 129291 : fd->have_reductemp = false;
228 129291 : fd->have_pointer_condtemp = false;
229 129291 : fd->have_scantemp = false;
230 129291 : fd->have_nonctrl_scantemp = false;
231 129291 : fd->non_rect = false;
232 129291 : fd->lastprivate_conditional = 0;
233 129291 : fd->tiling = NULL_TREE;
234 129291 : fd->collapse = 1;
235 129291 : fd->ordered = 0;
236 129291 : fd->first_nonrect = -1;
237 129291 : fd->last_nonrect = -1;
238 129291 : fd->sched_kind = OMP_CLAUSE_SCHEDULE_STATIC;
239 129291 : fd->sched_modifiers = 0;
240 129291 : fd->chunk_size = NULL_TREE;
241 129291 : fd->simd_schedule = false;
242 129291 : fd->first_inner_iterations = NULL_TREE;
243 129291 : fd->factor = NULL_TREE;
244 129291 : fd->adjn1 = NULL_TREE;
245 129291 : collapse_iter = NULL;
246 129291 : collapse_count = NULL;
247 :
248 647742 : for (t = gimple_omp_for_clauses (for_stmt); t ; t = OMP_CLAUSE_CHAIN (t))
249 518451 : switch (OMP_CLAUSE_CODE (t))
250 : {
251 45031 : case OMP_CLAUSE_NOWAIT:
252 45031 : fd->have_nowait = true;
253 45031 : break;
254 1399 : case OMP_CLAUSE_ORDERED:
255 1399 : fd->have_ordered = true;
256 1399 : if (OMP_CLAUSE_ORDERED_DOACROSS (t))
257 : {
258 700 : if (OMP_CLAUSE_ORDERED_EXPR (t))
259 680 : fd->ordered = tree_to_shwi (OMP_CLAUSE_ORDERED_EXPR (t));
260 : else
261 20 : fd->ordered = -1;
262 : }
263 : break;
264 30488 : case OMP_CLAUSE_SCHEDULE:
265 30488 : gcc_assert (!distribute && !taskloop);
266 30488 : fd->sched_kind
267 30488 : = (enum omp_clause_schedule_kind)
268 30488 : (OMP_CLAUSE_SCHEDULE_KIND (t) & OMP_CLAUSE_SCHEDULE_MASK);
269 30488 : fd->sched_modifiers = (OMP_CLAUSE_SCHEDULE_KIND (t)
270 30488 : & ~OMP_CLAUSE_SCHEDULE_MASK);
271 30488 : fd->chunk_size = OMP_CLAUSE_SCHEDULE_CHUNK_EXPR (t);
272 30488 : fd->simd_schedule = OMP_CLAUSE_SCHEDULE_SIMD (t);
273 30488 : break;
274 7026 : case OMP_CLAUSE_DIST_SCHEDULE:
275 7026 : gcc_assert (distribute);
276 7026 : fd->chunk_size = OMP_CLAUSE_DIST_SCHEDULE_CHUNK_EXPR (t);
277 7026 : break;
278 41507 : case OMP_CLAUSE_COLLAPSE:
279 41507 : fd->collapse = tree_to_shwi (OMP_CLAUSE_COLLAPSE_EXPR (t));
280 41507 : if (fd->collapse > 1)
281 : {
282 33914 : collapse_iter = &OMP_CLAUSE_COLLAPSE_ITERVAR (t);
283 33914 : collapse_count = &OMP_CLAUSE_COLLAPSE_COUNT (t);
284 : }
285 : break;
286 354 : case OMP_CLAUSE_TILE:
287 354 : fd->tiling = OMP_CLAUSE_TILE_LIST (t);
288 354 : fd->collapse = list_length (fd->tiling);
289 354 : gcc_assert (fd->collapse);
290 354 : collapse_iter = &OMP_CLAUSE_TILE_ITERVAR (t);
291 354 : collapse_count = &OMP_CLAUSE_TILE_COUNT (t);
292 354 : break;
293 482 : case OMP_CLAUSE__REDUCTEMP_:
294 482 : fd->have_reductemp = true;
295 482 : break;
296 40463 : case OMP_CLAUSE_LASTPRIVATE:
297 40463 : if (OMP_CLAUSE_LASTPRIVATE_CONDITIONAL (t))
298 1480 : fd->lastprivate_conditional++;
299 : break;
300 1526 : case OMP_CLAUSE__CONDTEMP_:
301 1526 : if (POINTER_TYPE_P (TREE_TYPE (OMP_CLAUSE_DECL (t))))
302 185 : fd->have_pointer_condtemp = true;
303 : break;
304 1126 : case OMP_CLAUSE__SCANTEMP_:
305 1126 : fd->have_scantemp = true;
306 1126 : if (!OMP_CLAUSE__SCANTEMP__ALLOC (t)
307 1126 : && !OMP_CLAUSE__SCANTEMP__CONTROL (t))
308 209 : fd->have_nonctrl_scantemp = true;
309 : break;
310 7718 : case OMP_CLAUSE_ORDER:
311 : /* FIXME: For OpenMP 5.2 this should change to
312 : if (OMP_CLAUSE_ORDER_REPRODUCIBLE (t))
313 : (with the exception of loop construct but that lowers to
314 : no schedule/dist_schedule clauses currently). */
315 7718 : if (!OMP_CLAUSE_ORDER_UNCONSTRAINED (t))
316 518451 : order_reproducible = true;
317 : default:
318 : break;
319 : }
320 :
321 129291 : if (fd->ordered == -1)
322 20 : fd->ordered = fd->collapse;
323 :
324 : /* For order(reproducible:concurrent) schedule ({dynamic,guided,runtime})
325 : we have either the option to expensively remember at runtime how we've
326 : distributed work from first loop and reuse that in following loops with
327 : the same number of iterations and schedule, or just force static schedule.
328 : OpenMP API calls etc. aren't allowed in order(concurrent) bodies so
329 : users can't observe it easily anyway. */
330 129291 : if (order_reproducible)
331 7423 : fd->sched_kind = OMP_CLAUSE_SCHEDULE_STATIC;
332 129291 : if (fd->collapse > 1 || fd->tiling)
333 34264 : fd->loops = loops;
334 : else
335 95027 : fd->loops = &fd->loop;
336 :
337 129291 : if (fd->ordered && fd->collapse == 1 && loops != NULL)
338 : {
339 188 : fd->loops = loops;
340 188 : iterv = NULL_TREE;
341 188 : countv = NULL_TREE;
342 188 : collapse_iter = &iterv;
343 188 : collapse_count = &countv;
344 : }
345 :
346 : /* FIXME: for now map schedule(auto) to schedule(static).
347 : There should be analysis to determine whether all iterations
348 : are approximately the same amount of work (then schedule(static)
349 : is best) or if it varies (then schedule(dynamic,N) is better). */
350 129291 : if (fd->sched_kind == OMP_CLAUSE_SCHEDULE_AUTO)
351 : {
352 5747 : fd->sched_kind = OMP_CLAUSE_SCHEDULE_STATIC;
353 5747 : gcc_assert (fd->chunk_size == NULL);
354 : }
355 129291 : gcc_assert ((fd->collapse == 1 && !fd->tiling) || collapse_iter != NULL);
356 129291 : if (taskloop)
357 9547 : fd->sched_kind = OMP_CLAUSE_SCHEDULE_RUNTIME;
358 129291 : if (fd->sched_kind == OMP_CLAUSE_SCHEDULE_RUNTIME)
359 16548 : gcc_assert (fd->chunk_size == NULL);
360 112743 : else if (fd->chunk_size == NULL)
361 : {
362 : /* We only need to compute a default chunk size for ordered
363 : static loops and dynamic loops. */
364 89612 : if (fd->sched_kind != OMP_CLAUSE_SCHEDULE_STATIC
365 89030 : || fd->have_ordered)
366 1303 : fd->chunk_size = (fd->sched_kind == OMP_CLAUSE_SCHEDULE_STATIC)
367 1303 : ? integer_zero_node : integer_one_node;
368 : }
369 :
370 129291 : int cnt = fd->ordered ? fd->ordered : fd->collapse;
371 129291 : int single_nonrect = -1;
372 129291 : tree single_nonrect_count = NULL_TREE;
373 129291 : enum tree_code single_nonrect_cond_code = ERROR_MARK;
374 193501 : for (i = 1; i < cnt; i++)
375 : {
376 64338 : tree n1 = gimple_omp_for_initial (for_stmt, i);
377 64338 : tree n2 = gimple_omp_for_final (for_stmt, i);
378 64338 : if (TREE_CODE (n1) == TREE_VEC)
379 : {
380 1786 : if (fd->non_rect)
381 : {
382 : single_nonrect = -1;
383 : break;
384 : }
385 2152 : for (int j = i - 1; j >= 0; j--)
386 2152 : if (TREE_VEC_ELT (n1, 0) == gimple_omp_for_index (for_stmt, j))
387 : {
388 : single_nonrect = j;
389 : break;
390 : }
391 1664 : fd->non_rect = true;
392 : }
393 62552 : else if (TREE_CODE (n2) == TREE_VEC)
394 : {
395 531 : if (fd->non_rect)
396 : {
397 : single_nonrect = -1;
398 : break;
399 : }
400 681 : for (int j = i - 1; j >= 0; j--)
401 681 : if (TREE_VEC_ELT (n2, 0) == gimple_omp_for_index (for_stmt, j))
402 : {
403 : single_nonrect = j;
404 : break;
405 : }
406 525 : fd->non_rect = true;
407 : }
408 : }
409 323004 : for (i = 0; i < cnt; i++)
410 : {
411 193713 : if (i == 0
412 129291 : && fd->collapse == 1
413 95199 : && !fd->tiling
414 95027 : && (fd->ordered == 0 || loops == NULL))
415 94839 : loop = &fd->loop;
416 98686 : else if (loops != NULL)
417 32768 : loop = loops + i;
418 : else
419 : loop = &dummy_loop;
420 :
421 193713 : loop->v = gimple_omp_for_index (for_stmt, i);
422 193713 : gcc_assert (SSA_VAR_P (loop->v));
423 193713 : gcc_assert (TREE_CODE (TREE_TYPE (loop->v)) == INTEGER_TYPE
424 : || TREE_CODE (TREE_TYPE (loop->v)) == BITINT_TYPE
425 : || TREE_CODE (TREE_TYPE (loop->v)) == POINTER_TYPE);
426 193713 : var = TREE_CODE (loop->v) == SSA_NAME ? SSA_NAME_VAR (loop->v) : loop->v;
427 193713 : loop->n1 = gimple_omp_for_initial (for_stmt, i);
428 193713 : loop->m1 = NULL_TREE;
429 193713 : loop->m2 = NULL_TREE;
430 193713 : loop->outer = 0;
431 193713 : loop->non_rect_referenced = false;
432 193713 : if (TREE_CODE (loop->n1) == TREE_VEC)
433 : {
434 2313 : for (int j = i - 1; j >= 0; j--)
435 2313 : if (TREE_VEC_ELT (loop->n1, 0) == gimple_omp_for_index (for_stmt, j))
436 : {
437 1786 : loop->outer = i - j;
438 1786 : if (loops != NULL)
439 639 : loops[j].non_rect_referenced = true;
440 1786 : if (fd->first_nonrect == -1 || fd->first_nonrect > j)
441 1664 : fd->first_nonrect = j;
442 : break;
443 : }
444 1786 : gcc_assert (loop->outer);
445 1786 : loop->m1 = TREE_VEC_ELT (loop->n1, 1);
446 1786 : loop->n1 = TREE_VEC_ELT (loop->n1, 2);
447 1786 : fd->non_rect = true;
448 1786 : fd->last_nonrect = i;
449 : }
450 :
451 193713 : loop->cond_code = gimple_omp_for_cond (for_stmt, i);
452 193713 : loop->n2 = gimple_omp_for_final (for_stmt, i);
453 193713 : gcc_assert (loop->cond_code != NE_EXPR
454 : || (gimple_omp_for_kind (for_stmt)
455 : != GF_OMP_FOR_KIND_OACC_LOOP));
456 193713 : if (TREE_CODE (loop->n2) == TREE_VEC)
457 : {
458 1273 : if (loop->outer)
459 742 : gcc_assert (TREE_VEC_ELT (loop->n2, 0)
460 : == gimple_omp_for_index (for_stmt, i - loop->outer));
461 : else
462 690 : for (int j = i - 1; j >= 0; j--)
463 690 : if (TREE_VEC_ELT (loop->n2, 0) == gimple_omp_for_index (for_stmt, j))
464 : {
465 531 : loop->outer = i - j;
466 531 : if (loops != NULL)
467 241 : loops[j].non_rect_referenced = true;
468 531 : if (fd->first_nonrect == -1 || fd->first_nonrect > j)
469 525 : fd->first_nonrect = j;
470 : break;
471 : }
472 1273 : gcc_assert (loop->outer);
473 1273 : loop->m2 = TREE_VEC_ELT (loop->n2, 1);
474 1273 : loop->n2 = TREE_VEC_ELT (loop->n2, 2);
475 1273 : fd->non_rect = true;
476 1273 : fd->last_nonrect = i;
477 : }
478 :
479 193713 : t = gimple_omp_for_incr (for_stmt, i);
480 193713 : gcc_assert (TREE_OPERAND (t, 0) == var);
481 193713 : loop->step = omp_get_for_step_from_incr (loc, t);
482 :
483 193713 : omp_adjust_for_condition (loc, &loop->cond_code, &loop->n2, loop->v,
484 : loop->step);
485 :
486 193713 : if (simd
487 153842 : || (fd->sched_kind == OMP_CLAUSE_SCHEDULE_STATIC
488 114457 : && !fd->have_ordered))
489 : {
490 152607 : if (fd->collapse == 1 && !fd->tiling)
491 77262 : iter_type = TREE_TYPE (loop->v);
492 75345 : else if (i == 0
493 75345 : || TYPE_PRECISION (iter_type)
494 48590 : < TYPE_PRECISION (TREE_TYPE (loop->v)))
495 : {
496 37632 : if (TREE_CODE (iter_type) == BITINT_TYPE
497 37632 : || TREE_CODE (TREE_TYPE (loop->v)) == BITINT_TYPE)
498 2 : iter_type
499 2 : = build_bitint_type (TYPE_PRECISION (TREE_TYPE (loop->v)),
500 : 1);
501 : else
502 37630 : iter_type
503 : = build_nonstandard_integer_type
504 37630 : (TYPE_PRECISION (TREE_TYPE (loop->v)), 1);
505 : }
506 : }
507 41106 : else if (iter_type != long_long_unsigned_type_node)
508 : {
509 37646 : if (POINTER_TYPE_P (TREE_TYPE (loop->v)))
510 : iter_type = long_long_unsigned_type_node;
511 35684 : else if (TYPE_UNSIGNED (TREE_TYPE (loop->v))
512 35684 : && TYPE_PRECISION (TREE_TYPE (loop->v))
513 6285 : >= TYPE_PRECISION (iter_type))
514 : {
515 2319 : tree n;
516 :
517 2319 : if (loop->cond_code == LT_EXPR)
518 451 : n = fold_build2_loc (loc, PLUS_EXPR, TREE_TYPE (loop->v),
519 : loop->n2, loop->step);
520 : else
521 1868 : n = loop->n1;
522 2319 : if (loop->m1
523 2319 : || loop->m2
524 2319 : || TREE_CODE (n) != INTEGER_CST
525 4284 : || tree_int_cst_lt (TYPE_MAX_VALUE (iter_type), n))
526 2242 : iter_type = long_long_unsigned_type_node;
527 : }
528 33365 : else if (TYPE_PRECISION (TREE_TYPE (loop->v))
529 33365 : > TYPE_PRECISION (iter_type))
530 : {
531 0 : tree n1, n2;
532 :
533 0 : if (loop->cond_code == LT_EXPR)
534 : {
535 0 : n1 = loop->n1;
536 0 : n2 = fold_build2_loc (loc, PLUS_EXPR, TREE_TYPE (loop->v),
537 : loop->n2, loop->step);
538 : }
539 : else
540 : {
541 0 : n1 = fold_build2_loc (loc, MINUS_EXPR, TREE_TYPE (loop->v),
542 : loop->n2, loop->step);
543 0 : n2 = loop->n1;
544 : }
545 0 : if (loop->m1
546 0 : || loop->m2
547 0 : || TREE_CODE (n1) != INTEGER_CST
548 0 : || TREE_CODE (n2) != INTEGER_CST
549 0 : || !tree_int_cst_lt (TYPE_MIN_VALUE (iter_type), n1)
550 0 : || !tree_int_cst_lt (n2, TYPE_MAX_VALUE (iter_type)))
551 0 : iter_type = long_long_unsigned_type_node;
552 : }
553 : }
554 :
555 193713 : if (i >= fd->collapse)
556 1802 : continue;
557 :
558 191911 : if (collapse_count && *collapse_count == NULL)
559 : {
560 31943 : if (count && integer_zerop (count))
561 2244 : continue;
562 29699 : tree n1first = NULL_TREE, n2first = NULL_TREE;
563 29699 : tree n1last = NULL_TREE, n2last = NULL_TREE;
564 29699 : tree ostep = NULL_TREE;
565 29699 : if (loop->m1 || loop->m2)
566 : {
567 911 : if (count == NULL_TREE)
568 645 : continue;
569 266 : if (single_nonrect == -1
570 247 : || (loop->m1 && TREE_CODE (loop->m1) != INTEGER_CST)
571 206 : || (loop->m2 && TREE_CODE (loop->m2) != INTEGER_CST)
572 190 : || TREE_CODE (loop->n1) != INTEGER_CST
573 186 : || TREE_CODE (loop->n2) != INTEGER_CST
574 182 : || TREE_CODE (loop->step) != INTEGER_CST)
575 : {
576 84 : count = NULL_TREE;
577 84 : continue;
578 : }
579 182 : tree var = gimple_omp_for_initial (for_stmt, single_nonrect);
580 182 : tree itype = TREE_TYPE (var);
581 182 : tree first = gimple_omp_for_initial (for_stmt, single_nonrect);
582 182 : t = gimple_omp_for_incr (for_stmt, single_nonrect);
583 182 : ostep = omp_get_for_step_from_incr (loc, t);
584 182 : t = fold_binary (MINUS_EXPR, long_long_unsigned_type_node,
585 : single_nonrect_count,
586 : build_one_cst (long_long_unsigned_type_node));
587 182 : t = fold_convert (itype, t);
588 182 : first = fold_convert (itype, first);
589 182 : ostep = fold_convert (itype, ostep);
590 182 : tree last = fold_binary (PLUS_EXPR, itype, first,
591 : fold_binary (MULT_EXPR, itype, t,
592 : ostep));
593 182 : if (TREE_CODE (first) != INTEGER_CST
594 182 : || TREE_CODE (last) != INTEGER_CST)
595 : {
596 0 : count = NULL_TREE;
597 0 : continue;
598 : }
599 182 : if (loop->m1)
600 : {
601 115 : tree m1 = fold_convert (itype, loop->m1);
602 115 : tree n1 = fold_convert (itype, loop->n1);
603 115 : n1first = fold_binary (PLUS_EXPR, itype,
604 : fold_binary (MULT_EXPR, itype,
605 : first, m1), n1);
606 115 : n1last = fold_binary (PLUS_EXPR, itype,
607 : fold_binary (MULT_EXPR, itype,
608 : last, m1), n1);
609 : }
610 : else
611 67 : n1first = n1last = loop->n1;
612 182 : if (loop->m2)
613 : {
614 136 : tree n2 = fold_convert (itype, loop->n2);
615 136 : tree m2 = fold_convert (itype, loop->m2);
616 136 : n2first = fold_binary (PLUS_EXPR, itype,
617 : fold_binary (MULT_EXPR, itype,
618 : first, m2), n2);
619 136 : n2last = fold_binary (PLUS_EXPR, itype,
620 : fold_binary (MULT_EXPR, itype,
621 : last, m2), n2);
622 : }
623 : else
624 46 : n2first = n2last = loop->n2;
625 182 : n1first = fold_convert (TREE_TYPE (loop->v), n1first);
626 182 : n2first = fold_convert (TREE_TYPE (loop->v), n2first);
627 182 : n1last = fold_convert (TREE_TYPE (loop->v), n1last);
628 182 : n2last = fold_convert (TREE_TYPE (loop->v), n2last);
629 182 : t = fold_binary (loop->cond_code, boolean_type_node,
630 : n1first, n2first);
631 182 : tree t2 = fold_binary (loop->cond_code, boolean_type_node,
632 : n1last, n2last);
633 182 : if (t && t2 && integer_nonzerop (t) && integer_nonzerop (t2))
634 : /* All outer loop iterators have at least one inner loop
635 : iteration. Try to compute the count at compile time. */
636 : t = NULL_TREE;
637 103 : else if (t && t2 && integer_zerop (t) && integer_zerop (t2))
638 : /* No iterations of the inner loop. count will be set to
639 : zero cst below. */;
640 95 : else if (TYPE_UNSIGNED (itype)
641 95 : || t == NULL_TREE
642 95 : || t2 == NULL_TREE
643 95 : || TREE_CODE (t) != INTEGER_CST
644 190 : || TREE_CODE (t2) != INTEGER_CST)
645 : {
646 : /* Punt (for now). */
647 0 : count = NULL_TREE;
648 0 : continue;
649 : }
650 : else
651 : {
652 : /* Some iterations of the outer loop have zero iterations
653 : of the inner loop, while others have at least one.
654 : In this case, we need to adjust one of those outer
655 : loop bounds. If ADJ_FIRST, we need to adjust outer n1
656 : (first), otherwise outer n2 (last). */
657 95 : bool adj_first = integer_zerop (t);
658 95 : tree n1 = fold_convert (itype, loop->n1);
659 95 : tree n2 = fold_convert (itype, loop->n2);
660 95 : tree m1 = loop->m1 ? fold_convert (itype, loop->m1)
661 25 : : build_zero_cst (itype);
662 95 : tree m2 = loop->m2 ? fold_convert (itype, loop->m2)
663 11 : : build_zero_cst (itype);
664 95 : t = fold_binary (MINUS_EXPR, itype, n1, n2);
665 95 : t2 = fold_binary (MINUS_EXPR, itype, m2, m1);
666 95 : t = fold_binary (TRUNC_DIV_EXPR, itype, t, t2);
667 95 : t2 = fold_binary (MINUS_EXPR, itype, t, first);
668 95 : t2 = fold_binary (TRUNC_MOD_EXPR, itype, t2, ostep);
669 95 : t = fold_binary (MINUS_EXPR, itype, t, t2);
670 95 : tree n1cur
671 95 : = fold_binary (PLUS_EXPR, itype, n1,
672 : fold_binary (MULT_EXPR, itype, m1, t));
673 95 : tree n2cur
674 95 : = fold_binary (PLUS_EXPR, itype, n2,
675 : fold_binary (MULT_EXPR, itype, m2, t));
676 95 : t2 = fold_binary (loop->cond_code, boolean_type_node,
677 : n1cur, n2cur);
678 95 : tree t3 = fold_binary (MULT_EXPR, itype, m1, ostep);
679 95 : tree t4 = fold_binary (MULT_EXPR, itype, m2, ostep);
680 95 : tree diff;
681 95 : if (adj_first)
682 : {
683 67 : tree new_first;
684 67 : if (integer_nonzerop (t2))
685 : {
686 6 : new_first = t;
687 6 : n1first = n1cur;
688 6 : n2first = n2cur;
689 6 : if (flag_checking)
690 : {
691 6 : t3 = fold_binary (MINUS_EXPR, itype, n1cur, t3);
692 6 : t4 = fold_binary (MINUS_EXPR, itype, n2cur, t4);
693 6 : t3 = fold_binary (loop->cond_code,
694 : boolean_type_node, t3, t4);
695 6 : gcc_assert (integer_zerop (t3));
696 : }
697 : }
698 : else
699 : {
700 61 : t3 = fold_binary (PLUS_EXPR, itype, n1cur, t3);
701 61 : t4 = fold_binary (PLUS_EXPR, itype, n2cur, t4);
702 61 : new_first = fold_binary (PLUS_EXPR, itype, t, ostep);
703 61 : n1first = t3;
704 61 : n2first = t4;
705 61 : if (flag_checking)
706 : {
707 61 : t3 = fold_binary (loop->cond_code,
708 : boolean_type_node, t3, t4);
709 61 : gcc_assert (integer_nonzerop (t3));
710 : }
711 : }
712 67 : diff = fold_binary (MINUS_EXPR, itype, new_first, first);
713 67 : first = new_first;
714 67 : fd->adjn1 = first;
715 : }
716 : else
717 : {
718 28 : tree new_last;
719 28 : if (integer_zerop (t2))
720 : {
721 11 : t3 = fold_binary (MINUS_EXPR, itype, n1cur, t3);
722 11 : t4 = fold_binary (MINUS_EXPR, itype, n2cur, t4);
723 11 : new_last = fold_binary (MINUS_EXPR, itype, t, ostep);
724 11 : n1last = t3;
725 11 : n2last = t4;
726 11 : if (flag_checking)
727 : {
728 11 : t3 = fold_binary (loop->cond_code,
729 : boolean_type_node, t3, t4);
730 11 : gcc_assert (integer_nonzerop (t3));
731 : }
732 : }
733 : else
734 : {
735 17 : new_last = t;
736 17 : n1last = n1cur;
737 17 : n2last = n2cur;
738 17 : if (flag_checking)
739 : {
740 17 : t3 = fold_binary (PLUS_EXPR, itype, n1cur, t3);
741 17 : t4 = fold_binary (PLUS_EXPR, itype, n2cur, t4);
742 17 : t3 = fold_binary (loop->cond_code,
743 : boolean_type_node, t3, t4);
744 17 : gcc_assert (integer_zerop (t3));
745 : }
746 : }
747 28 : diff = fold_binary (MINUS_EXPR, itype, last, new_last);
748 : }
749 95 : if (TYPE_UNSIGNED (itype)
750 95 : && single_nonrect_cond_code == GT_EXPR)
751 0 : diff = fold_binary (TRUNC_DIV_EXPR, itype,
752 : fold_unary (NEGATE_EXPR, itype, diff),
753 : fold_unary (NEGATE_EXPR, itype,
754 : ostep));
755 : else
756 95 : diff = fold_binary (TRUNC_DIV_EXPR, itype, diff, ostep);
757 95 : diff = fold_convert (long_long_unsigned_type_node, diff);
758 95 : single_nonrect_count
759 95 : = fold_binary (MINUS_EXPR, long_long_unsigned_type_node,
760 : single_nonrect_count, diff);
761 95 : t = NULL_TREE;
762 : }
763 : }
764 : else
765 28788 : t = fold_binary (loop->cond_code, boolean_type_node,
766 : fold_convert (TREE_TYPE (loop->v), loop->n1),
767 : fold_convert (TREE_TYPE (loop->v), loop->n2));
768 28891 : if (t && integer_zerop (t))
769 2252 : count = build_zero_cst (long_long_unsigned_type_node);
770 26718 : else if ((i == 0 || count != NULL_TREE)
771 17002 : && (TREE_CODE (TREE_TYPE (loop->v)) == INTEGER_TYPE
772 1202 : || TREE_CODE (TREE_TYPE (loop->v)) == BITINT_TYPE)
773 15801 : && TREE_CONSTANT (loop->n1)
774 11797 : && TREE_CONSTANT (loop->n2)
775 37508 : && TREE_CODE (loop->step) == INTEGER_CST)
776 : {
777 10768 : tree itype = TREE_TYPE (loop->v);
778 :
779 10768 : if (POINTER_TYPE_P (itype))
780 0 : itype = signed_type_for (itype);
781 13136 : t = build_int_cst (itype, (loop->cond_code == LT_EXPR ? -1 : 1));
782 10768 : t = fold_build2 (PLUS_EXPR, itype,
783 : fold_convert (itype, loop->step), t);
784 10768 : tree n1 = loop->n1;
785 10768 : tree n2 = loop->n2;
786 10768 : if (loop->m1 || loop->m2)
787 : {
788 174 : gcc_assert (single_nonrect != -1);
789 : n1 = n1first;
790 : n2 = n2first;
791 : }
792 10768 : t = fold_build2 (PLUS_EXPR, itype, t, fold_convert (itype, n2));
793 10768 : t = fold_build2 (MINUS_EXPR, itype, t, fold_convert (itype, n1));
794 10768 : tree step = fold_convert_loc (loc, itype, loop->step);
795 10768 : if (TYPE_UNSIGNED (itype) && loop->cond_code == GT_EXPR)
796 2254 : t = fold_build2 (TRUNC_DIV_EXPR, itype,
797 : fold_build1 (NEGATE_EXPR, itype, t),
798 : fold_build1 (NEGATE_EXPR, itype, step));
799 : else
800 8514 : t = fold_build2 (TRUNC_DIV_EXPR, itype, t, step);
801 10768 : tree llutype = long_long_unsigned_type_node;
802 10768 : t = fold_convert (llutype, t);
803 10768 : if (loop->m1 || loop->m2)
804 : {
805 : /* t is number of iterations of inner loop at either first
806 : or last value of the outer iterator (the one with fewer
807 : iterations).
808 : Compute t2 = ((m2 - m1) * ostep) / step
809 : and niters = outer_count * t
810 : + t2 * ((outer_count - 1) * outer_count / 2)
811 : */
812 174 : tree m1 = loop->m1 ? loop->m1 : integer_zero_node;
813 174 : tree m2 = loop->m2 ? loop->m2 : integer_zero_node;
814 174 : m1 = fold_convert (itype, m1);
815 174 : m2 = fold_convert (itype, m2);
816 174 : tree t2 = fold_build2 (MINUS_EXPR, itype, m2, m1);
817 174 : t2 = fold_build2 (MULT_EXPR, itype, t2, ostep);
818 174 : if (TYPE_UNSIGNED (itype) && loop->cond_code == GT_EXPR)
819 0 : t2 = fold_build2 (TRUNC_DIV_EXPR, itype,
820 : fold_build1 (NEGATE_EXPR, itype, t2),
821 : fold_build1 (NEGATE_EXPR, itype, step));
822 : else
823 174 : t2 = fold_build2 (TRUNC_DIV_EXPR, itype, t2, step);
824 174 : t2 = fold_convert (llutype, t2);
825 174 : fd->first_inner_iterations = t;
826 174 : fd->factor = t2;
827 174 : t = fold_build2 (MULT_EXPR, llutype, t,
828 : single_nonrect_count);
829 174 : tree t3 = fold_build2 (MINUS_EXPR, llutype,
830 : single_nonrect_count,
831 : build_one_cst (llutype));
832 174 : t3 = fold_build2 (MULT_EXPR, llutype, t3,
833 : single_nonrect_count);
834 174 : t3 = fold_build2 (TRUNC_DIV_EXPR, llutype, t3,
835 : build_int_cst (llutype, 2));
836 174 : t2 = fold_build2 (MULT_EXPR, llutype, t2, t3);
837 174 : t = fold_build2 (PLUS_EXPR, llutype, t, t2);
838 : }
839 10768 : if (i == single_nonrect)
840 : {
841 247 : if (integer_zerop (t) || TREE_CODE (t) != INTEGER_CST)
842 : count = t;
843 : else
844 : {
845 247 : single_nonrect_count = t;
846 247 : single_nonrect_cond_code = loop->cond_code;
847 247 : if (count == NULL_TREE)
848 246 : count = build_one_cst (llutype);
849 : }
850 : }
851 10521 : else if (count != NULL_TREE)
852 4207 : count = fold_build2 (MULT_EXPR, llutype, count, t);
853 : else
854 : count = t;
855 10768 : if (TREE_CODE (count) != INTEGER_CST)
856 0 : count = NULL_TREE;
857 : }
858 15950 : else if (count && !integer_zerop (count))
859 : count = NULL_TREE;
860 : }
861 : }
862 :
863 129291 : if (count
864 129291 : && !simd
865 4050 : && (fd->sched_kind != OMP_CLAUSE_SCHEDULE_STATIC
866 3297 : || fd->have_ordered))
867 : {
868 853 : if (!tree_int_cst_lt (count, TYPE_MAX_VALUE (long_integer_type_node)))
869 0 : iter_type = long_long_unsigned_type_node;
870 : else
871 853 : iter_type = long_integer_type_node;
872 : }
873 128438 : else if (collapse_iter && *collapse_iter != NULL)
874 22907 : iter_type = TREE_TYPE (*collapse_iter);
875 129291 : fd->iter_type = iter_type;
876 129291 : if (collapse_iter && *collapse_iter == NULL)
877 11549 : *collapse_iter = create_tmp_var (iter_type, ".iter");
878 129291 : if (collapse_count && *collapse_count == NULL)
879 : {
880 11549 : if (count)
881 : {
882 5231 : *collapse_count = fold_convert_loc (loc, iter_type, count);
883 5231 : if (fd->first_inner_iterations && fd->factor)
884 : {
885 174 : t = make_tree_vec (4);
886 174 : TREE_VEC_ELT (t, 0) = *collapse_count;
887 174 : TREE_VEC_ELT (t, 1) = fd->first_inner_iterations;
888 174 : TREE_VEC_ELT (t, 2) = fd->factor;
889 174 : TREE_VEC_ELT (t, 3) = fd->adjn1;
890 174 : *collapse_count = t;
891 : }
892 : }
893 : else
894 6318 : *collapse_count = create_tmp_var (iter_type, ".count");
895 : }
896 :
897 129291 : if (fd->collapse > 1 || fd->tiling || (fd->ordered && loops))
898 : {
899 34452 : fd->loop.v = *collapse_iter;
900 34452 : fd->loop.n1 = build_int_cst (TREE_TYPE (fd->loop.v), 0);
901 34452 : fd->loop.n2 = *collapse_count;
902 34452 : if (TREE_CODE (fd->loop.n2) == TREE_VEC)
903 : {
904 377 : gcc_assert (fd->non_rect);
905 377 : fd->first_inner_iterations = TREE_VEC_ELT (fd->loop.n2, 1);
906 377 : fd->factor = TREE_VEC_ELT (fd->loop.n2, 2);
907 377 : fd->adjn1 = TREE_VEC_ELT (fd->loop.n2, 3);
908 377 : fd->loop.n2 = TREE_VEC_ELT (fd->loop.n2, 0);
909 : }
910 34452 : fd->loop.step = build_int_cst (TREE_TYPE (fd->loop.v), 1);
911 34452 : fd->loop.m1 = NULL_TREE;
912 34452 : fd->loop.m2 = NULL_TREE;
913 34452 : fd->loop.outer = 0;
914 34452 : fd->loop.cond_code = LT_EXPR;
915 : }
916 94623 : else if (loops)
917 36074 : loops[0] = fd->loop;
918 129291 : }
919 :
920 : /* Build a call to GOMP_barrier. */
921 :
922 : gimple *
923 4701 : omp_build_barrier (tree lhs, int kind)
924 : {
925 9368 : tree fndecl = builtin_decl_explicit (lhs ? BUILT_IN_GOMP_BARRIER_CANCEL
926 : : BUILT_IN_GOMP_BARRIER);
927 4701 : gcall *g = gimple_build_call (fndecl, 1,
928 4701 : build_int_cst (integer_type_node, kind));
929 4701 : if (lhs)
930 34 : gimple_call_set_lhs (g, lhs);
931 4701 : return g;
932 : }
933 :
934 : /* Find OMP_FOR resp. OMP_SIMD with non-NULL OMP_FOR_INIT. Also, fill in pdata
935 : array, pdata[0] non-NULL if there is anything non-trivial in between,
936 : pdata[1] is address of OMP_PARALLEL in between if any, pdata[2] is address
937 : of OMP_FOR in between if any and pdata[3] is address of the inner
938 : OMP_FOR/OMP_SIMD. */
939 :
940 : tree
941 110126 : find_combined_omp_for (tree *tp, int *walk_subtrees, void *data)
942 : {
943 110126 : tree **pdata = (tree **) data;
944 110126 : *walk_subtrees = 0;
945 110126 : switch (TREE_CODE (*tp))
946 : {
947 11876 : case OMP_FOR:
948 11876 : if (OMP_FOR_INIT (*tp) != NULL_TREE)
949 : {
950 6057 : pdata[3] = tp;
951 6057 : return *tp;
952 : }
953 5819 : pdata[2] = tp;
954 5819 : *walk_subtrees = 1;
955 5819 : break;
956 15322 : case OMP_SIMD:
957 15322 : if (OMP_FOR_INIT (*tp) != NULL_TREE)
958 : {
959 15322 : pdata[3] = tp;
960 15322 : return *tp;
961 : }
962 : break;
963 46745 : case BIND_EXPR:
964 46745 : if (BIND_EXPR_VARS (*tp)
965 46745 : || (BIND_EXPR_BLOCK (*tp)
966 41141 : && BLOCK_VARS (BIND_EXPR_BLOCK (*tp))))
967 5149 : pdata[0] = tp;
968 46745 : *walk_subtrees = 1;
969 46745 : break;
970 8573 : case STATEMENT_LIST:
971 8573 : if (!tsi_one_before_end_p (tsi_start (*tp)))
972 241 : pdata[0] = tp;
973 8573 : *walk_subtrees = 1;
974 8573 : break;
975 216 : case TRY_FINALLY_EXPR:
976 216 : case CLEANUP_POINT_EXPR:
977 216 : pdata[0] = tp;
978 216 : *walk_subtrees = 1;
979 216 : break;
980 11904 : case OMP_PARALLEL:
981 11904 : pdata[1] = tp;
982 11904 : *walk_subtrees = 1;
983 11904 : break;
984 : default:
985 : break;
986 : }
987 : return NULL_TREE;
988 : }
989 :
990 : /* Return maximum possible vectorization factor for the target, or for
991 : the OpenMP offload target if one exists. */
992 :
993 : poly_uint64
994 33481 : omp_max_vf (bool offload)
995 : {
996 33481 : if (!optimize
997 32143 : || optimize_debug
998 32143 : || !flag_tree_loop_optimize
999 32142 : || (!flag_tree_loop_vectorize
1000 585 : && OPTION_SET_P (flag_tree_loop_vectorize)))
1001 1343 : return 1;
1002 :
1003 32138 : if (ENABLE_OFFLOADING && offload)
1004 : {
1005 : for (const char *c = getenv ("OFFLOAD_TARGET_NAMES"); c;)
1006 : {
1007 : if (startswith (c, "amdgcn"))
1008 : return ordered_max (poly_uint64 (64), omp_max_vf (false));
1009 : else if ((c = strchr (c, ':')))
1010 : c++;
1011 : }
1012 : /* Otherwise, fall through to host VF. */
1013 : }
1014 :
1015 32138 : auto_vector_modes modes;
1016 32138 : targetm.vectorize.autovectorize_vector_modes (&modes, true);
1017 32138 : if (!modes.is_empty ())
1018 : {
1019 : poly_uint64 vf = 0;
1020 134018 : for (unsigned int i = 0; i < modes.length (); ++i)
1021 : /* The returned modes use the smallest element size (and thus
1022 : the largest nunits) for the vectorization approach that they
1023 : represent. */
1024 203768 : vf = ordered_max (vf, GET_MODE_NUNITS (modes[i]));
1025 32134 : return vf;
1026 : }
1027 :
1028 4 : machine_mode vqimode = targetm.vectorize.preferred_simd_mode (QImode);
1029 4 : if (GET_MODE_CLASS (vqimode) == MODE_VECTOR_INT)
1030 0 : return GET_MODE_NUNITS (vqimode);
1031 :
1032 4 : return 1;
1033 32138 : }
1034 :
1035 : /* Return maximum SIMT width if offloading may target SIMT hardware. */
1036 :
1037 : int
1038 3641 : omp_max_simt_vf (void)
1039 : {
1040 3641 : if (!optimize)
1041 : return 0;
1042 : if (ENABLE_OFFLOADING)
1043 : for (const char *c = getenv ("OFFLOAD_TARGET_NAMES"); c;)
1044 : {
1045 : if (startswith (c, "nvptx"))
1046 : return 32;
1047 : else if ((c = strchr (c, ':')))
1048 : c++;
1049 : }
1050 : return 0;
1051 : }
1052 :
1053 : /* Return true if PROP is possibly present in one of the offloading target's
1054 : OpenMP contexts. The format of PROPS string is always offloading target's
1055 : name terminated by '\0', followed by properties for that offloading
1056 : target separated by '\0' and terminated by another '\0'. The strings
1057 : are created from omp-device-properties installed files of all configured
1058 : offloading targets. */
1059 :
1060 : static bool
1061 0 : omp_offload_device_kind_arch_isa (const char *props, const char *prop)
1062 : {
1063 0 : const char *names = getenv ("OFFLOAD_TARGET_NAMES");
1064 0 : if (names == NULL || *names == '\0')
1065 : return false;
1066 0 : while (*props != '\0')
1067 : {
1068 0 : size_t name_len = strlen (props);
1069 0 : bool matches = false;
1070 0 : for (const char *c = names; c; )
1071 : {
1072 0 : if (strncmp (props, c, name_len) == 0
1073 0 : && (c[name_len] == '\0'
1074 : || c[name_len] == ':'
1075 : || c[name_len] == '='))
1076 : {
1077 : matches = true;
1078 : break;
1079 : }
1080 0 : else if ((c = strchr (c, ':')))
1081 0 : c++;
1082 : }
1083 0 : props = props + name_len + 1;
1084 0 : while (*props != '\0')
1085 : {
1086 0 : if (matches && strcmp (props, prop) == 0)
1087 : return true;
1088 0 : props = strchr (props, '\0') + 1;
1089 : }
1090 0 : props++;
1091 : }
1092 : return false;
1093 : }
1094 :
1095 : /* Return true if the current code location is or might be offloaded.
1096 : Return true in declare target functions, or when nested in a target
1097 : region or when unsure, return false otherwise. */
1098 :
1099 : static bool
1100 0 : omp_maybe_offloaded (tree construct_context)
1101 : {
1102 : /* No offload targets available? */
1103 0 : if (!ENABLE_OFFLOADING)
1104 0 : return false;
1105 : const char *names = getenv ("OFFLOAD_TARGET_NAMES");
1106 : if (names == NULL || *names == '\0')
1107 : return false;
1108 :
1109 : /* Parsing is too early to tell. */
1110 : if (symtab->state == PARSING)
1111 : /* Maybe. */
1112 : return true;
1113 :
1114 : /* Late resolution of offloaded code happens in the offload compiler,
1115 : where it's treated as native code instead. So return false here. */
1116 : if (cfun && cfun->after_inlining)
1117 : return false;
1118 :
1119 : /* Check if the function is marked for offloading (either explicitly
1120 : or via omp_discover_implicit_declare_target). */
1121 : if (current_function_decl
1122 : && lookup_attribute ("omp declare target",
1123 : DECL_ATTRIBUTES (current_function_decl)))
1124 : return true;
1125 :
1126 : /* Check for nesting inside a target directive. */
1127 : for (tree ts = construct_context; ts; ts = TREE_CHAIN (ts))
1128 : if (OMP_TS_CODE (ts) == OMP_TRAIT_CONSTRUCT_TARGET)
1129 : return true;
1130 :
1131 : return false;
1132 : }
1133 :
1134 : /* Lookup tables for context selectors. */
1135 : const char *omp_tss_map[] =
1136 : {
1137 : "construct",
1138 : "device",
1139 : "target_device",
1140 : "implementation",
1141 : "user",
1142 : NULL
1143 : };
1144 :
1145 : /* Arrays of property candidates must be null-terminated. */
1146 : static const char *const kind_properties[] =
1147 : { "host", "nohost", "cpu", "gpu", "fpga", "any", NULL };
1148 : static const char *const vendor_properties[] =
1149 : { "amd", "arm", "bsc", "cray", "fujitsu", "gnu", "hpe", "ibm", "intel",
1150 : "llvm", "nec", "nvidia", "pgi", "ti", "unknown", NULL };
1151 : static const char *const extension_properties[] =
1152 : { NULL };
1153 : static const char *const atomic_default_mem_order_properties[] =
1154 : { "seq_cst", "relaxed", "acq_rel", "acquire", "release", NULL };
1155 :
1156 : struct omp_ts_info omp_ts_map[] =
1157 : {
1158 : { "kind",
1159 : (1 << OMP_TRAIT_SET_DEVICE) | (1 << OMP_TRAIT_SET_TARGET_DEVICE),
1160 : OMP_TRAIT_PROPERTY_NAME_LIST, false,
1161 : kind_properties
1162 : },
1163 : { "isa",
1164 : (1 << OMP_TRAIT_SET_DEVICE) | (1 << OMP_TRAIT_SET_TARGET_DEVICE),
1165 : OMP_TRAIT_PROPERTY_NAME_LIST, false,
1166 : NULL
1167 : },
1168 : { "arch",
1169 : (1 << OMP_TRAIT_SET_DEVICE) | (1 << OMP_TRAIT_SET_TARGET_DEVICE),
1170 : OMP_TRAIT_PROPERTY_NAME_LIST, false,
1171 : NULL
1172 : },
1173 : { "device_num",
1174 : (1 << OMP_TRAIT_SET_TARGET_DEVICE),
1175 : OMP_TRAIT_PROPERTY_DEV_NUM_EXPR, false,
1176 : NULL
1177 : },
1178 : { "vendor",
1179 : (1 << OMP_TRAIT_SET_IMPLEMENTATION),
1180 : OMP_TRAIT_PROPERTY_NAME_LIST, true,
1181 : vendor_properties,
1182 : },
1183 : { "extension",
1184 : (1 << OMP_TRAIT_SET_IMPLEMENTATION),
1185 : OMP_TRAIT_PROPERTY_NAME_LIST, true,
1186 : extension_properties,
1187 : },
1188 : { "atomic_default_mem_order",
1189 : (1 << OMP_TRAIT_SET_IMPLEMENTATION),
1190 : OMP_TRAIT_PROPERTY_ID, true,
1191 : atomic_default_mem_order_properties,
1192 : },
1193 : { "requires",
1194 : (1 << OMP_TRAIT_SET_IMPLEMENTATION),
1195 : OMP_TRAIT_PROPERTY_CLAUSE_LIST, true,
1196 : NULL
1197 : },
1198 : { "unified_address",
1199 : (1 << OMP_TRAIT_SET_IMPLEMENTATION),
1200 : OMP_TRAIT_PROPERTY_NONE, true,
1201 : NULL
1202 : },
1203 : { "unified_shared_memory",
1204 : (1 << OMP_TRAIT_SET_IMPLEMENTATION),
1205 : OMP_TRAIT_PROPERTY_NONE, true,
1206 : NULL
1207 : },
1208 : { "self_maps",
1209 : (1 << OMP_TRAIT_SET_IMPLEMENTATION),
1210 : OMP_TRAIT_PROPERTY_NONE, true,
1211 : NULL
1212 : },
1213 : { "dynamic_allocators",
1214 : (1 << OMP_TRAIT_SET_IMPLEMENTATION),
1215 : OMP_TRAIT_PROPERTY_NONE, true,
1216 : NULL
1217 : },
1218 : { "reverse_offload",
1219 : (1 << OMP_TRAIT_SET_IMPLEMENTATION),
1220 : OMP_TRAIT_PROPERTY_NONE, true,
1221 : NULL
1222 : },
1223 : { "condition",
1224 : (1 << OMP_TRAIT_SET_USER),
1225 : OMP_TRAIT_PROPERTY_BOOL_EXPR, true,
1226 : NULL
1227 : },
1228 : { "target",
1229 : (1 << OMP_TRAIT_SET_CONSTRUCT),
1230 : OMP_TRAIT_PROPERTY_NONE, false,
1231 : NULL
1232 : },
1233 : { "teams",
1234 : (1 << OMP_TRAIT_SET_CONSTRUCT),
1235 : OMP_TRAIT_PROPERTY_NONE, false,
1236 : NULL
1237 : },
1238 : { "parallel",
1239 : (1 << OMP_TRAIT_SET_CONSTRUCT),
1240 : OMP_TRAIT_PROPERTY_NONE, false,
1241 : NULL
1242 : },
1243 : { "for",
1244 : (1 << OMP_TRAIT_SET_CONSTRUCT),
1245 : OMP_TRAIT_PROPERTY_NONE, false,
1246 : NULL
1247 : },
1248 : { "simd",
1249 : (1 << OMP_TRAIT_SET_CONSTRUCT),
1250 : OMP_TRAIT_PROPERTY_CLAUSE_LIST, false,
1251 : NULL
1252 : },
1253 : { "dispatch",
1254 : (1 << OMP_TRAIT_SET_CONSTRUCT),
1255 : OMP_TRAIT_PROPERTY_NONE, false,
1256 : NULL
1257 : },
1258 : { NULL, 0, OMP_TRAIT_PROPERTY_NONE, false, NULL } /* OMP_TRAIT_LAST */
1259 : };
1260 :
1261 : /* Return a name from PROP, a property in selectors accepting
1262 : name lists. */
1263 :
1264 : const char *
1265 13748 : omp_context_name_list_prop (tree prop)
1266 : {
1267 13748 : gcc_assert (OMP_TP_NAME (prop) == OMP_TP_NAMELIST_NODE);
1268 13748 : tree val = OMP_TP_VALUE (prop);
1269 13748 : switch (TREE_CODE (val))
1270 : {
1271 8351 : case IDENTIFIER_NODE:
1272 8351 : return IDENTIFIER_POINTER (val);
1273 5397 : case STRING_CST:
1274 : #ifdef ACCEL_COMPILER
1275 : return TREE_STRING_POINTER (val);
1276 : #else
1277 5397 : {
1278 5397 : const char *ret = TREE_STRING_POINTER (val);
1279 10794 : if ((size_t) TREE_STRING_LENGTH (val)
1280 10318 : == strlen (ret) + (lang_GNU_Fortran () ? 0 : 1))
1281 : return ret;
1282 : return NULL;
1283 : }
1284 : #endif
1285 : default:
1286 : return NULL;
1287 : }
1288 : }
1289 :
1290 :
1291 : /* Helper function called via walk_tree, to determine if *TP is a
1292 : PARM_DECL. */
1293 : static tree
1294 1437 : expr_uses_parm_decl (tree *tp, int *walk_subtrees ATTRIBUTE_UNUSED,
1295 : void *data ATTRIBUTE_UNUSED)
1296 : {
1297 1437 : if (TREE_CODE (*tp) == PARM_DECL)
1298 26 : return *tp;
1299 : return NULL_TREE;
1300 : }
1301 :
1302 : /* Diagnose errors in an OpenMP context selector, return CTX if
1303 : it is correct or error_mark_node otherwise. */
1304 :
1305 : tree
1306 3686 : omp_check_context_selector (location_t loc, tree ctx,
1307 : enum omp_ctx_directive directive)
1308 : {
1309 3686 : bool tss_seen[OMP_TRAIT_SET_LAST], ts_seen[OMP_TRAIT_LAST];
1310 :
1311 3686 : memset (tss_seen, 0, sizeof (tss_seen));
1312 7844 : for (tree tss = ctx; tss; tss = TREE_CHAIN (tss))
1313 : {
1314 4288 : enum omp_tss_code tss_code = OMP_TSS_CODE (tss);
1315 4288 : bool saw_any_prop = false;
1316 4288 : bool saw_other_prop = false;
1317 :
1318 : /* Each trait-set-selector-name can only be specified once. */
1319 4288 : if (tss_seen[tss_code])
1320 : {
1321 60 : error_at (loc, "selector set %qs specified more than once",
1322 30 : OMP_TSS_NAME (tss));
1323 30 : return error_mark_node;
1324 : }
1325 : else
1326 4258 : tss_seen[tss_code] = true;
1327 :
1328 4258 : memset (ts_seen, 0, sizeof (ts_seen));
1329 9307 : for (tree ts = OMP_TSS_TRAIT_SELECTORS (tss); ts; ts = TREE_CHAIN (ts))
1330 : {
1331 5129 : enum omp_ts_code ts_code = OMP_TS_CODE (ts);
1332 :
1333 : /* Ignore unknown traits. */
1334 5129 : if (ts_code == OMP_TRAIT_INVALID)
1335 73 : continue;
1336 :
1337 : /* Each trait-selector-name can only be specified once. */
1338 5056 : if (ts_seen[ts_code])
1339 : {
1340 15 : error_at (loc,
1341 : "selector %qs specified more than once in set %qs",
1342 15 : OMP_TS_NAME (ts),
1343 15 : OMP_TSS_NAME (tss));
1344 15 : return error_mark_node;
1345 : }
1346 : else
1347 5041 : ts_seen[ts_code] = true;
1348 :
1349 : /* If trait-property "any" is specified in the "kind"
1350 : trait-selector of the "device" selector set or the
1351 : "target_device" selector sets, no other trait-property
1352 : may be specified in the same selector set. */
1353 5041 : if (ts_code == OMP_TRAIT_DEVICE_KIND)
1354 1306 : for (tree p = OMP_TS_PROPERTIES (ts); p; p = TREE_CHAIN (p))
1355 : {
1356 478 : const char *prop = omp_context_name_list_prop (p);
1357 478 : if (!prop)
1358 4 : continue;
1359 474 : else if (strcmp (prop, "any") == 0)
1360 : saw_any_prop = true;
1361 : else
1362 399 : saw_other_prop = true;
1363 : }
1364 : /* It seems slightly suspicious that the spec's language covers
1365 : the device_num selector too, but
1366 : target_device={device_num(whatever),kind(any)}
1367 : is probably not terribly useful anyway. */
1368 4627 : else if (ts_code == OMP_TRAIT_DEVICE_ARCH
1369 : || ts_code == OMP_TRAIT_DEVICE_ISA
1370 4627 : || ts_code == OMP_TRAIT_DEVICE_NUM)
1371 1052 : saw_other_prop = true;
1372 :
1373 : /* Each trait-property can only be specified once in a trait-selector
1374 : other than the construct selector set. FIXME: only handles
1375 : name-list properties, not clause-list properties, since the
1376 : "requires" selector is not implemented yet (PR 113067). */
1377 5041 : if (tss_code != OMP_TRAIT_SET_CONSTRUCT)
1378 8292 : for (tree p1 = OMP_TS_PROPERTIES (ts); p1; p1 = TREE_CHAIN (p1))
1379 : {
1380 3199 : if (OMP_TP_NAME (p1) != OMP_TP_NAMELIST_NODE)
1381 : break;
1382 2268 : const char *n1 = omp_context_name_list_prop (p1);
1383 2268 : if (!n1)
1384 8 : continue;
1385 2618 : for (tree p2 = TREE_CHAIN (p1); p2; p2 = TREE_CHAIN (p2))
1386 : {
1387 388 : const char *n2 = omp_context_name_list_prop (p2);
1388 388 : if (!n2)
1389 0 : continue;
1390 388 : if (!strcmp (n1, n2))
1391 : {
1392 30 : error_at (loc,
1393 : "trait-property %qs specified more "
1394 : "than once in %qs selector",
1395 30 : n1, OMP_TS_NAME (ts));
1396 30 : return error_mark_node;
1397 : }
1398 : }
1399 : }
1400 :
1401 : /* This restriction is documented in the spec in the section
1402 : for the metadirective "when" clause (7.4.1 in the 5.2 spec). */
1403 5011 : if (directive == OMP_CTX_METADIRECTIVE
1404 5011 : && ts_code == OMP_TRAIT_CONSTRUCT_SIMD
1405 5015 : && OMP_TS_PROPERTIES (ts))
1406 : {
1407 0 : error_at (loc,
1408 : "properties must not be specified for the %<simd%> "
1409 : "selector in a %<metadirective%> context-selector");
1410 0 : return error_mark_node;
1411 : }
1412 :
1413 : /* "simd" is not allowed at all in "begin declare variant"
1414 : selectors. */
1415 5011 : if (directive == OMP_CTX_BEGIN_DECLARE_VARIANT
1416 5011 : && ts_code == OMP_TRAIT_CONSTRUCT_SIMD)
1417 : {
1418 4 : error_at (loc,
1419 : "the %<simd%> selector is not permitted in a "
1420 : "%<begin declare variant%> context selector");
1421 4 : return error_mark_node;
1422 : }
1423 :
1424 : /* Reject expressions that reference parameter variables in
1425 : "declare variant", as this is not yet implemented. FIXME;
1426 : see PR middle-end/113904. */
1427 5007 : if (directive != OMP_CTX_METADIRECTIVE
1428 4135 : && (ts_code == OMP_TRAIT_DEVICE_NUM
1429 4135 : || ts_code == OMP_TRAIT_USER_CONDITION))
1430 : {
1431 552 : tree exp = OMP_TS_PROPERTIES (ts);
1432 552 : if (walk_tree (&exp, expr_uses_parm_decl, NULL, NULL))
1433 : {
1434 26 : sorry_at (loc,
1435 : "reference to function parameter in "
1436 : "%<declare variant%> dynamic selector expression");
1437 26 : return error_mark_node;
1438 : }
1439 : }
1440 :
1441 : /* Check for unknown properties. */
1442 4981 : if (omp_ts_map[ts_code].valid_properties == NULL)
1443 3799 : continue;
1444 3630 : for (tree p = OMP_TS_PROPERTIES (ts); p; p = TREE_CHAIN (p))
1445 4959 : for (unsigned j = 0; ; j++)
1446 : {
1447 6230 : const char *candidate
1448 6230 : = omp_ts_map[ts_code].valid_properties[j];
1449 6230 : if (candidate == NULL)
1450 : {
1451 : /* We've reached the end of the candidate array. */
1452 83 : if (ts_code == OMP_TRAIT_IMPLEMENTATION_ADMO)
1453 : /* FIXME: not sure why this is an error vs warnings
1454 : for the others, + incorrect/unknown wording? */
1455 : {
1456 5 : error_at (loc,
1457 : "incorrect property %qs of %qs selector",
1458 5 : IDENTIFIER_POINTER (OMP_TP_NAME (p)),
1459 : "atomic_default_mem_order");
1460 5 : return error_mark_node;
1461 : }
1462 78 : if (OMP_TP_NAME (p) == OMP_TP_NAMELIST_NODE
1463 78 : && (TREE_CODE (OMP_TP_VALUE (p)) == STRING_CST))
1464 46 : warning_at (loc, OPT_Wopenmp,
1465 : "unknown property %qE of %qs selector",
1466 23 : OMP_TP_VALUE (p),
1467 23 : OMP_TS_NAME (ts));
1468 55 : else if (OMP_TP_NAME (p) == OMP_TP_NAMELIST_NODE)
1469 110 : warning_at (loc, OPT_Wopenmp,
1470 : "unknown property %qs of %qs selector",
1471 : omp_context_name_list_prop (p),
1472 55 : OMP_TS_NAME (ts));
1473 0 : else if (OMP_TP_NAME (p))
1474 0 : warning_at (loc, OPT_Wopenmp,
1475 : "unknown property %qs of %qs selector",
1476 0 : IDENTIFIER_POINTER (OMP_TP_NAME (p)),
1477 0 : OMP_TS_NAME (ts));
1478 : break;
1479 : }
1480 6147 : else if (OMP_TP_NAME (p) == OMP_TP_NAMELIST_NODE)
1481 : /* Property-list traits. */
1482 : {
1483 5899 : const char *str = omp_context_name_list_prop (p);
1484 5899 : if (str && !strcmp (str, candidate))
1485 : break;
1486 : }
1487 248 : else if (!strcmp (IDENTIFIER_POINTER (OMP_TP_NAME (p)),
1488 : candidate))
1489 : /* Identifier traits. */
1490 : break;
1491 4959 : }
1492 : }
1493 :
1494 4178 : if (saw_any_prop && saw_other_prop)
1495 : {
1496 20 : error_at (loc,
1497 : "no other trait-property may be specified "
1498 : "in the same selector set with %<kind(\"any\")%>");
1499 20 : return error_mark_node;
1500 : }
1501 : }
1502 : return ctx;
1503 : }
1504 :
1505 : /* Produce a mangled version of BASE_ID for the name of the variant
1506 : function with context selector CTX. SEP is a separator string.
1507 : The return value is an IDENTIFIER_NODE.
1508 :
1509 : Per the OpenMP spec, "the symbol names of two definitions of a function are
1510 : considered to be equal if and only if their effective context selectors are
1511 : equivalent". However, if we did have two such definitions, we'd get an ODR
1512 : violation. We already take steps in the front ends to make variant
1513 : functions internal to the compilation unit, since there is no (portable) way
1514 : to reference them directly by name or declare them as extern in another
1515 : compilation unit. So, we can diagnose the would-be ODR violations by
1516 : checking that there is not already a variant for the same function with an
1517 : equivalent context selector, and otherwise just use a simple counter to name
1518 : the variant functions instead of any complicated scheme to encode the
1519 : context selector in the name.
1520 :
1521 : C++ and Fortran modules are an exception to this, as variants in a module
1522 : interface unit are visible to implementation TUs for that module. This
1523 : is handled in the C++ front end by adding an additional prefix to SEP for
1524 : variants in a module interface to prevent collisions with names in the
1525 : other TUs. The Fortran and C++ front ends add the module name to the
1526 : output of this function using their normal mechanisms for symbols with
1527 : module linkage. */
1528 : tree
1529 308 : omp_mangle_variant_name (tree base_id, tree ctx ATTRIBUTE_UNUSED,
1530 : const char *sep)
1531 : {
1532 308 : const char *base_name = IDENTIFIER_POINTER (base_id);
1533 :
1534 : /* Now do the actual mangling. */
1535 308 : static int variant_counter;
1536 : /* The numeric suffix and terminating byte ought to need way less than
1537 : 32 bytes extra, that's just a magic number. */
1538 308 : size_t buflen = (strlen (base_name) + strlen (sep) + strlen ("ompvariant")
1539 : + 32);
1540 308 : char *buffer = (char *) alloca (buflen);
1541 308 : buflen = snprintf (buffer, buflen, "%s%sompvariant%d", base_name, sep,
1542 : ++variant_counter);
1543 308 : return get_identifier_with_length (buffer, buflen);
1544 : }
1545 :
1546 : /* Forward declaration. */
1547 : static int omp_context_selector_compare (tree ctx1, tree ctx2);
1548 :
1549 : /* Diagnose an error if there is already a variant with CTX registered
1550 : for BASE_DECL. Returns true if OK, false otherwise. */
1551 : bool
1552 233 : omp_check_for_duplicate_variant (location_t loc, tree base_decl, tree ctx)
1553 : {
1554 262 : for (tree attr = DECL_ATTRIBUTES (base_decl); attr; attr = TREE_CHAIN (attr))
1555 : {
1556 33 : attr = lookup_attribute ("omp declare variant base", attr);
1557 33 : if (attr == NULL_TREE)
1558 : break;
1559 :
1560 33 : tree selector = TREE_VALUE (TREE_VALUE (attr));
1561 33 : if (omp_context_selector_compare (ctx, selector) == 0)
1562 : {
1563 4 : error_at (loc,
1564 : "multiple definitions of variants with the same "
1565 : "context selector violate the one-definition rule");
1566 4 : inform (DECL_SOURCE_LOCATION (TREE_PURPOSE (TREE_VALUE (attr))),
1567 : "previous variant declaration here");
1568 4 : return false;
1569 : }
1570 : }
1571 : return true;
1572 : }
1573 :
1574 : /* Forward declarations. */
1575 : static int omp_context_selector_set_compare (enum omp_tss_code, tree, tree);
1576 : static int omp_construct_simd_compare (tree, tree, bool);
1577 :
1578 : /* Register VARIANT as variant of some base function marked with
1579 : #pragma omp declare variant. CONSTRUCT is corresponding list of
1580 : trait-selectors for the construct selector set. This is stashed as the
1581 : value of the "omp declare variant variant" attribute on VARIANT. */
1582 : void
1583 2642 : omp_mark_declare_variant (location_t loc, tree variant, tree construct)
1584 : {
1585 : /* Ignore this variant if it contains unknown construct selectors.
1586 : It will never match, and the front ends have already issued a warning
1587 : about it. */
1588 4246 : for (tree c = construct; c; c = TREE_CHAIN (c))
1589 1660 : if (OMP_TS_CODE (c) == OMP_TRAIT_INVALID)
1590 : return;
1591 :
1592 2586 : tree attr = lookup_attribute ("omp declare variant variant",
1593 2586 : DECL_ATTRIBUTES (variant));
1594 2586 : if (attr == NULL_TREE)
1595 : {
1596 1882 : attr = tree_cons (get_identifier ("omp declare variant variant"),
1597 : unshare_expr (construct),
1598 1882 : DECL_ATTRIBUTES (variant));
1599 1882 : DECL_ATTRIBUTES (variant) = attr;
1600 1882 : return;
1601 : }
1602 704 : if ((TREE_VALUE (attr) != NULL_TREE) != (construct != NULL_TREE)
1603 704 : || (construct != NULL_TREE
1604 154 : && omp_context_selector_set_compare (OMP_TRAIT_SET_CONSTRUCT,
1605 154 : TREE_VALUE (attr),
1606 : construct)))
1607 53 : error_at (loc, "%qD used as a variant with incompatible %<construct%> "
1608 : "selector sets", variant);
1609 : }
1610 :
1611 :
1612 : /* Constructors for context selectors. */
1613 :
1614 : tree
1615 4363 : make_trait_set_selector (enum omp_tss_code code, tree selectors, tree chain)
1616 : {
1617 4363 : return tree_cons (build_int_cst (integer_type_node, code),
1618 4363 : selectors, chain);
1619 : }
1620 :
1621 : tree
1622 7701 : make_trait_selector (enum omp_ts_code code, tree score, tree properties,
1623 : tree chain)
1624 : {
1625 7701 : if (score == NULL_TREE)
1626 7260 : return tree_cons (build_int_cst (integer_type_node, code),
1627 : properties, chain);
1628 : else
1629 441 : return tree_cons (build_int_cst (integer_type_node, code),
1630 : tree_cons (OMP_TS_SCORE_NODE, score, properties),
1631 : chain);
1632 : }
1633 :
1634 : tree
1635 3318 : make_trait_property (tree name, tree value, tree chain)
1636 : {
1637 3318 : return tree_cons (name, value, chain);
1638 : }
1639 :
1640 : /* Constructor for metadirective variants. */
1641 : tree
1642 904 : make_omp_metadirective_variant (tree selector, tree directive, tree body)
1643 : {
1644 904 : return build_tree_list (selector, build_tree_list (directive, body));
1645 : }
1646 :
1647 : /* If the construct selector traits SELECTOR_TRAITS match the corresponding
1648 : OpenMP context traits CONTEXT_TRAITS, return true and set *SCORE to the
1649 : corresponding score if it is non-null. */
1650 : static bool
1651 4699 : omp_construct_traits_match (tree selector_traits, tree context_traits,
1652 : score_wide_int *score)
1653 : {
1654 4699 : int slength = list_length (selector_traits);
1655 4699 : int clength = list_length (context_traits);
1656 :
1657 : /* Trivial failure: the selector has more traits than the OpenMP context. */
1658 4699 : if (slength > clength)
1659 : return false;
1660 :
1661 : /* There's only one trait in the selector and it doesn't have any properties
1662 : to match. */
1663 8072 : if (slength == 1 && !OMP_TS_PROPERTIES (selector_traits))
1664 : {
1665 3786 : int p = 0, i = 1;
1666 3786 : enum omp_ts_code code = OMP_TS_CODE (selector_traits);
1667 7868 : for (tree t = context_traits; t; t = TREE_CHAIN (t), i++)
1668 4082 : if (OMP_TS_CODE (t) == code)
1669 3819 : p = i;
1670 3786 : if (p != 0)
1671 : {
1672 3759 : if (score)
1673 1317 : *score = wi::shifted_mask <score_wide_int> (p - 1, 1, false);
1674 3759 : return true;
1675 : }
1676 : else
1677 : return false;
1678 : }
1679 :
1680 : /* Now handle the more general cases.
1681 : Both lists of traits are ordered from outside in, corresponding to
1682 : the c1, ..., cN numbering for the OpenMP context specified in
1683 : in section 7.1 of the OpenMP 5.2 spec. Section 7.3 of the spec says
1684 : "if the traits that correspond to the construct selector set appear
1685 : multiple times in the OpenMP context, the highest valued subset of
1686 : context traits that contains all trait selectors in the same order
1687 : are used". This means that we want to start the search for a match
1688 : from the end of the list, rather than the beginning. To facilitate
1689 : that, transfer the lists to temporary arrays to allow random access
1690 : to the elements (their order remains outside in). */
1691 492 : int i, j;
1692 492 : tree s, c;
1693 :
1694 492 : tree *sarray = (tree *) alloca (slength * sizeof (tree));
1695 1574 : for (s = selector_traits, i = 0; s; s = TREE_CHAIN (s), i++)
1696 1082 : sarray[i] = s;
1697 :
1698 492 : tree *carray = (tree *) alloca (clength * sizeof (tree));
1699 1785 : for (c = context_traits, j = 0; c; c = TREE_CHAIN (c), j++)
1700 1293 : carray[j] = c;
1701 :
1702 : /* The variable "i" indexes the selector, "j" indexes the OpenMP context.
1703 : Find the "j" corresponding to each sarray[i]. Note that the spec uses
1704 : "p" as the 1-based position, but "j" is zero-based, e.g. equal to
1705 : p - 1. */
1706 492 : score_wide_int result = 0;
1707 492 : j = clength - 1;
1708 1452 : for (i = slength - 1; i >= 0; i--)
1709 : {
1710 1029 : enum omp_ts_code code = OMP_TS_CODE (sarray[i]);
1711 1029 : tree props = OMP_TS_PROPERTIES (sarray[i]);
1712 1287 : for (; j >= 0; j--)
1713 : {
1714 1218 : if (OMP_TS_CODE (carray[j]) != code)
1715 218 : continue;
1716 1040 : if (code == OMP_TRAIT_CONSTRUCT_SIMD
1717 1000 : && props
1718 1040 : && omp_construct_simd_compare (props,
1719 40 : OMP_TS_PROPERTIES (carray[j]),
1720 : true) > 0)
1721 40 : continue;
1722 : break;
1723 : }
1724 : /* If j >= 0, we have a match for this trait at position j. */
1725 960 : if (j < 0)
1726 : return false;
1727 960 : result += wi::shifted_mask <score_wide_int> (j, 1, false);
1728 960 : j--;
1729 : }
1730 423 : if (score)
1731 141 : *score = result;
1732 : return true;
1733 : }
1734 :
1735 : /* Return 1 if context selector CTX matches the current OpenMP context, 0
1736 : if it does not and -1 if it is unknown and need to be determined later.
1737 : Some properties can be checked right away during parsing, others need
1738 : to wait until the whole TU is parsed, others need to wait until
1739 : IPA, others until vectorization.
1740 :
1741 : CONSTRUCT_CONTEXT is a list of construct traits from the OpenMP context,
1742 : which must be collected by omp_get_construct_context during
1743 : gimplification. It is ignored (and may be null) if this function is
1744 : called during parsing. Otherwise COMPLETE_P should indicate whether
1745 : CONSTRUCT_CONTEXT is known to be complete and not missing constructs
1746 : filled in later during compilation.
1747 :
1748 : If DECLARE_VARIANT_ELISION_P is true, the function implements the test
1749 : for elision of preprocessed code in "begin declare variant" constructs,
1750 : and returns 0 only for failure to match traits in the device and
1751 : implementation sets.
1752 :
1753 : Dynamic properties (which are evaluated at run-time) should always
1754 : return 1. */
1755 :
1756 : int
1757 10745 : omp_context_selector_matches (tree ctx,
1758 : tree construct_context,
1759 : bool complete_p,
1760 : bool declare_variant_elision_p)
1761 : {
1762 10745 : int ret = 1;
1763 10745 : bool maybe_offloaded = omp_maybe_offloaded (construct_context);
1764 :
1765 20542 : for (tree tss = ctx; tss; tss = TREE_CHAIN (tss))
1766 : {
1767 11536 : enum omp_tss_code set = OMP_TSS_CODE (tss);
1768 11536 : tree selectors = OMP_TSS_TRAIT_SELECTORS (tss);
1769 :
1770 : /* Immediately reject the match if there are any ignored
1771 : selectors present. */
1772 11536 : if (!declare_variant_elision_p
1773 11536 : || set == OMP_TRAIT_SET_DEVICE
1774 369 : || set == OMP_TRAIT_SET_IMPLEMENTATION)
1775 24776 : for (tree ts = selectors; ts; ts = TREE_CHAIN (ts))
1776 13387 : if (OMP_TS_CODE (ts) == OMP_TRAIT_INVALID)
1777 : return 0;
1778 :
1779 11468 : if (set == OMP_TRAIT_SET_CONSTRUCT)
1780 : {
1781 : /* We cannot resolve the construct selector during parsing because
1782 : the OpenMP context (and CONSTRUCT_CONTEXT) isn't available
1783 : until gimplification. */
1784 4794 : if (symtab->state == PARSING)
1785 : {
1786 1557 : ret = -1;
1787 1557 : continue;
1788 : }
1789 :
1790 3237 : gcc_assert (selectors);
1791 :
1792 : /* During gimplification, CONSTRUCT_CONTEXT is partial, and doesn't
1793 : include a construct for "declare simd" that may be added
1794 : when there is not an enclosing "target" construct. We might
1795 : be able to find a positive match against the partial context
1796 : (although we cannot yet score it accurately), but if we can't,
1797 : treat it as unknown instead of no match. */
1798 3237 : if (!omp_construct_traits_match (selectors, construct_context, NULL))
1799 : {
1800 : /* If we've got a complete context, it's definitely a failed
1801 : match. */
1802 513 : if (complete_p)
1803 : return 0;
1804 :
1805 : /* If the selector doesn't include simd, then we don't have
1806 : to worry about whether "declare simd" would cause it to
1807 : match; so this is also a definite failure. */
1808 8 : bool have_simd = false;
1809 8 : for (tree ts = selectors; ts; ts = TREE_CHAIN (ts))
1810 8 : if (OMP_TS_CODE (ts) == OMP_TRAIT_CONSTRUCT_SIMD)
1811 : {
1812 : have_simd = true;
1813 : break;
1814 : }
1815 8 : if (!have_simd)
1816 : return 0;
1817 : else
1818 : ret = -1;
1819 : }
1820 2732 : continue;
1821 2732 : }
1822 6674 : else if (set == OMP_TRAIT_SET_TARGET_DEVICE)
1823 : /* The target_device set is dynamic, so treat it as always
1824 : resolvable. However, the current implementation doesn't
1825 : support it in a target region, so diagnose that as an error.
1826 : FIXME: maybe make this a warning and return 0 instead? */
1827 : {
1828 489 : for (tree ts = construct_context; ts; ts = TREE_CHAIN (ts))
1829 0 : if (OMP_TS_CODE (ts) == OMP_TRAIT_CONSTRUCT_TARGET)
1830 0 : sorry ("%<target_device%> selector set inside of %<target%> "
1831 : "directive");
1832 489 : continue;
1833 489 : }
1834 :
1835 11499 : for (tree ts = selectors; ts; ts = TREE_CHAIN (ts))
1836 : {
1837 6480 : enum omp_ts_code sel = OMP_TS_CODE (ts);
1838 6480 : switch (sel)
1839 : {
1840 1591 : case OMP_TRAIT_IMPLEMENTATION_VENDOR:
1841 1591 : gcc_assert (set == OMP_TRAIT_SET_IMPLEMENTATION);
1842 4618 : for (tree p = OMP_TS_PROPERTIES (ts); p; p = TREE_CHAIN (p))
1843 : {
1844 1606 : const char *prop = omp_context_name_list_prop (p);
1845 1606 : if (prop == NULL)
1846 : return 0;
1847 1602 : if (!strcmp (prop, "gnu"))
1848 1436 : continue;
1849 : return 0;
1850 : }
1851 : break;
1852 30 : case OMP_TRAIT_IMPLEMENTATION_EXTENSION:
1853 30 : gcc_assert (set == OMP_TRAIT_SET_IMPLEMENTATION);
1854 : /* We don't support any extensions right now. */
1855 : return 0;
1856 163 : break;
1857 163 : case OMP_TRAIT_IMPLEMENTATION_ADMO:
1858 163 : gcc_assert (set == OMP_TRAIT_SET_IMPLEMENTATION);
1859 163 : if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
1860 : break;
1861 :
1862 163 : {
1863 163 : enum omp_memory_order omo
1864 : = ((enum omp_memory_order)
1865 163 : (omp_requires_mask
1866 : & OMP_REQUIRES_ATOMIC_DEFAULT_MEM_ORDER));
1867 163 : if (omo == OMP_MEMORY_ORDER_UNSPECIFIED)
1868 : {
1869 : /* We don't know yet, until end of TU. */
1870 72 : if (symtab->state == PARSING)
1871 : {
1872 : ret = -1;
1873 : break;
1874 : }
1875 : else
1876 : omo = OMP_MEMORY_ORDER_RELAXED;
1877 : }
1878 91 : tree p = OMP_TS_PROPERTIES (ts);
1879 91 : const char *prop = IDENTIFIER_POINTER (OMP_TP_NAME (p));
1880 91 : if (!strcmp (prop, "relaxed")
1881 5 : && omo != OMP_MEMORY_ORDER_RELAXED)
1882 : return 0;
1883 86 : else if (!strcmp (prop, "seq_cst")
1884 86 : && omo != OMP_MEMORY_ORDER_SEQ_CST)
1885 : return 0;
1886 86 : else if (!strcmp (prop, "acq_rel")
1887 0 : && omo != OMP_MEMORY_ORDER_ACQ_REL)
1888 : return 0;
1889 86 : else if (!strcmp (prop, "acquire")
1890 0 : && omo != OMP_MEMORY_ORDER_ACQUIRE)
1891 : return 0;
1892 86 : else if (!strcmp (prop, "release")
1893 0 : && omo != OMP_MEMORY_ORDER_RELEASE)
1894 : return 0;
1895 : }
1896 : break;
1897 734 : case OMP_TRAIT_DEVICE_ARCH:
1898 734 : gcc_assert (set == OMP_TRAIT_SET_DEVICE);
1899 1655 : for (tree p = OMP_TS_PROPERTIES (ts); p; p = TREE_CHAIN (p))
1900 : {
1901 739 : const char *arch = omp_context_name_list_prop (p);
1902 739 : if (arch == NULL)
1903 : return 0;
1904 739 : int r = 0;
1905 739 : if (targetm.omp.device_kind_arch_isa != NULL)
1906 739 : r = targetm.omp.device_kind_arch_isa (omp_device_arch,
1907 : arch);
1908 739 : if (r == 0 || (r == -1 && symtab->state != PARSING))
1909 : {
1910 : /* If we are or might be in a target region or
1911 : declare target function, need to take into account
1912 : also offloading values.
1913 : Note that maybe_offloaded is always false in late
1914 : resolution; that's handled as native code (the
1915 : above case) in the offload compiler instead. */
1916 552 : if (!maybe_offloaded)
1917 552 : return 0;
1918 : if (ENABLE_OFFLOADING)
1919 : {
1920 : const char *arches = omp_offload_device_arch;
1921 : if (omp_offload_device_kind_arch_isa (arches, arch))
1922 : {
1923 : ret = -1;
1924 : continue;
1925 : }
1926 : }
1927 : return 0;
1928 : }
1929 : else if (r == -1)
1930 : ret = -1;
1931 : /* If arch matches on the host, it still might not match
1932 : in the offloading region. */
1933 : else if (maybe_offloaded)
1934 : ret = -1;
1935 : }
1936 : break;
1937 20 : case OMP_TRAIT_IMPLEMENTATION_UNIFIED_ADDRESS:
1938 20 : gcc_assert (set == OMP_TRAIT_SET_IMPLEMENTATION);
1939 20 : if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
1940 : break;
1941 :
1942 20 : if ((omp_requires_mask & OMP_REQUIRES_UNIFIED_ADDRESS) == 0)
1943 : {
1944 20 : if (symtab->state == PARSING)
1945 : ret = -1;
1946 : else
1947 : return 0;
1948 : }
1949 : break;
1950 15 : case OMP_TRAIT_IMPLEMENTATION_UNIFIED_SHARED_MEMORY:
1951 15 : gcc_assert (set == OMP_TRAIT_SET_IMPLEMENTATION);
1952 15 : if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
1953 : break;
1954 :
1955 15 : if ((omp_requires_mask
1956 : & OMP_REQUIRES_UNIFIED_SHARED_MEMORY) == 0)
1957 : {
1958 15 : if (symtab->state == PARSING)
1959 : ret = -1;
1960 : else
1961 : return 0;
1962 : }
1963 : break;
1964 5 : case OMP_TRAIT_IMPLEMENTATION_SELF_MAPS:
1965 5 : gcc_assert (set == OMP_TRAIT_SET_IMPLEMENTATION);
1966 5 : if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
1967 : break;
1968 :
1969 5 : if ((omp_requires_mask & OMP_REQUIRES_SELF_MAPS) == 0)
1970 : {
1971 5 : if (symtab->state == PARSING)
1972 : ret = -1;
1973 : else
1974 : return 0;
1975 : }
1976 : break;
1977 10 : case OMP_TRAIT_IMPLEMENTATION_DYNAMIC_ALLOCATORS:
1978 10 : gcc_assert (set == OMP_TRAIT_SET_IMPLEMENTATION);
1979 10 : if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
1980 : break;
1981 :
1982 10 : if ((omp_requires_mask
1983 : & OMP_REQUIRES_DYNAMIC_ALLOCATORS) == 0)
1984 : {
1985 10 : if (symtab->state == PARSING)
1986 : ret = -1;
1987 : else
1988 : return 0;
1989 : }
1990 : break;
1991 10 : case OMP_TRAIT_IMPLEMENTATION_REVERSE_OFFLOAD:
1992 10 : gcc_assert (set == OMP_TRAIT_SET_IMPLEMENTATION);
1993 10 : if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
1994 : break;
1995 :
1996 10 : if ((omp_requires_mask & OMP_REQUIRES_REVERSE_OFFLOAD) == 0)
1997 : {
1998 10 : if (symtab->state == PARSING)
1999 : ret = -1;
2000 : else
2001 : return 0;
2002 : }
2003 : break;
2004 714 : case OMP_TRAIT_DEVICE_KIND:
2005 714 : gcc_assert (set == OMP_TRAIT_SET_DEVICE);
2006 2084 : for (tree p = OMP_TS_PROPERTIES (ts); p; p = TREE_CHAIN (p))
2007 : {
2008 738 : const char *prop = omp_context_name_list_prop (p);
2009 738 : if (prop == NULL)
2010 : return 0;
2011 734 : if (!strcmp (prop, "any"))
2012 42 : continue;
2013 692 : if (!strcmp (prop, "host"))
2014 : {
2015 : #ifdef ACCEL_COMPILER
2016 : return 0;
2017 : #else
2018 345 : if (maybe_offloaded)
2019 : ret = -1;
2020 345 : continue;
2021 : #endif
2022 : }
2023 347 : if (!strcmp (prop, "nohost"))
2024 : {
2025 : #ifndef ACCEL_COMPILER
2026 : if (maybe_offloaded)
2027 : ret = -1;
2028 : else
2029 : return 0;
2030 : #endif
2031 : continue;
2032 : }
2033 :
2034 319 : int r = 0;
2035 319 : if (targetm.omp.device_kind_arch_isa != NULL)
2036 319 : r = targetm.omp.device_kind_arch_isa (omp_device_kind,
2037 : prop);
2038 : else
2039 : #ifndef ACCEL_COMPILER
2040 0 : r = strcmp (prop, "cpu") == 0;
2041 : #else
2042 : gcc_unreachable ();
2043 : #endif
2044 319 : if (r == 0 || (r == -1 && symtab->state != PARSING))
2045 : {
2046 : /* If we are or might be in a target region or
2047 : declare target function, need to take into account
2048 : also offloading values.
2049 : Note that maybe_offloaded is always false in late
2050 : resolution; that's handled as native code (the
2051 : above case) in the offload compiler instead. */
2052 : if (!maybe_offloaded)
2053 : return 0;
2054 : if (ENABLE_OFFLOADING)
2055 : {
2056 : const char *kinds = omp_offload_device_kind;
2057 : if (omp_offload_device_kind_arch_isa (kinds, prop))
2058 : {
2059 : ret = -1;
2060 : continue;
2061 : }
2062 : }
2063 : return 0;
2064 : }
2065 : else if (r == -1)
2066 : ret = -1;
2067 : /* If kind matches on the host, it still might not match
2068 : in the offloading region. */
2069 : else if (maybe_offloaded)
2070 : ret = -1;
2071 : }
2072 : break;
2073 543 : case OMP_TRAIT_DEVICE_ISA:
2074 543 : gcc_assert (set == OMP_TRAIT_SET_DEVICE);
2075 1732 : for (tree p = OMP_TS_PROPERTIES (ts); p; p = TREE_CHAIN (p))
2076 : {
2077 801 : const char *isa = omp_context_name_list_prop (p);
2078 801 : if (isa == NULL)
2079 : return 0;
2080 801 : int r = 0;
2081 801 : if (targetm.omp.device_kind_arch_isa != NULL)
2082 801 : r = targetm.omp.device_kind_arch_isa (omp_device_isa,
2083 : isa);
2084 801 : if (r == 0 || (r == -1 && symtab->state != PARSING))
2085 : {
2086 : /* If isa is valid on the target, but not in the
2087 : current function and current function has
2088 : #pragma omp declare simd on it, some simd clones
2089 : might have the isa added later on. */
2090 139 : if (r == -1
2091 139 : && targetm.simd_clone.compute_vecsize_and_simdlen
2092 139 : && (cfun == NULL || !cfun->after_inlining))
2093 : {
2094 55 : tree attrs
2095 55 : = DECL_ATTRIBUTES (current_function_decl);
2096 55 : if (lookup_attribute ("omp declare simd", attrs))
2097 : {
2098 34 : ret = -1;
2099 34 : continue;
2100 : }
2101 : }
2102 : /* If we are or might be in a target region or
2103 : declare target function, need to take into account
2104 : also offloading values.
2105 : Note that maybe_offloaded is always false in late
2106 : resolution; that's handled as native code (the
2107 : above case) in the offload compiler instead. */
2108 155 : if (!maybe_offloaded)
2109 155 : return 0;
2110 : if (ENABLE_OFFLOADING)
2111 : {
2112 : const char *isas = omp_offload_device_isa;
2113 : if (omp_offload_device_kind_arch_isa (isas, isa))
2114 : {
2115 : ret = -1;
2116 : continue;
2117 : }
2118 : }
2119 : return 0;
2120 : }
2121 : else if (r == -1)
2122 : ret = -1;
2123 : /* If isa matches on the host, it still might not match
2124 : in the offloading region. */
2125 : else if (maybe_offloaded)
2126 : ret = -1;
2127 : }
2128 : break;
2129 2645 : case OMP_TRAIT_USER_CONDITION:
2130 2645 : gcc_assert (set == OMP_TRAIT_SET_USER);
2131 : /* The spec does not include the "user" set in the things that
2132 : can trigger code elision in "begin declare variant". */
2133 2645 : if (declare_variant_elision_p)
2134 : {
2135 : ret = -1;
2136 : break;
2137 : }
2138 5266 : for (tree p = OMP_TS_PROPERTIES (ts); p; p = TREE_CHAIN (p))
2139 2633 : if (OMP_TP_NAME (p) == NULL_TREE)
2140 : {
2141 : /* If the expression is not a constant, the selector
2142 : is dynamic. */
2143 2633 : if (!tree_fits_shwi_p (OMP_TP_VALUE (p)))
2144 : break;
2145 :
2146 1001 : if (integer_zerop (OMP_TP_VALUE (p)))
2147 : return 0;
2148 829 : if (integer_nonzerop (OMP_TP_VALUE (p)))
2149 : break;
2150 : ret = -1;
2151 : }
2152 : break;
2153 : case OMP_TRAIT_INVALID:
2154 : /* This is only for the declare_variant_elision_p case. */
2155 5314 : ret = -1;
2156 : break;
2157 : default:
2158 : break;
2159 : }
2160 : }
2161 : }
2162 : return ret;
2163 : }
2164 :
2165 : /* Helper function for resolve_omp_target_device_matches, also used
2166 : directly when we know in advance that the device is the host to avoid
2167 : the overhead of late resolution. SEL is the selector code and
2168 : PROPERTIES are the properties to match. The return value is a
2169 : boolean. */
2170 : static bool
2171 153 : omp_target_device_matches_on_host (enum omp_ts_code selector,
2172 : tree properties)
2173 : {
2174 153 : bool result = 1;
2175 :
2176 153 : if (dump_file)
2177 0 : fprintf (dump_file, "omp_target_device_matches_on_host:\n");
2178 :
2179 153 : switch (selector)
2180 : {
2181 : case OMP_TRAIT_DEVICE_KIND:
2182 216 : for (tree p = properties; p && result; p = TREE_CHAIN (p))
2183 : {
2184 108 : const char *prop = omp_context_name_list_prop (p);
2185 :
2186 108 : if (prop == NULL)
2187 : result = 0;
2188 108 : else if (!strcmp (prop, "any"))
2189 : ;
2190 108 : else if (!strcmp (prop, "host"))
2191 : {
2192 : #ifdef ACCEL_COMPILER
2193 : result = 0;
2194 : #else
2195 : ;
2196 : #endif
2197 : }
2198 52 : else if (!strcmp (prop, "nohost"))
2199 : {
2200 : #ifdef ACCEL_COMPILER
2201 : ;
2202 : #else
2203 : result = 0;
2204 : #endif
2205 : }
2206 44 : else if (targetm.omp.device_kind_arch_isa != NULL)
2207 44 : result = targetm.omp.device_kind_arch_isa (omp_device_kind, prop);
2208 : else
2209 : #ifndef ACCEL_COMPILER
2210 0 : result = strcmp (prop, "cpu") == 0;
2211 : #else
2212 : gcc_unreachable ();
2213 : #endif
2214 108 : if (dump_file)
2215 0 : fprintf (dump_file, "Matching device kind %s = %s\n",
2216 : prop, (result ? "true" : "false"));
2217 : }
2218 : break;
2219 29 : case OMP_TRAIT_DEVICE_ARCH:
2220 29 : if (targetm.omp.device_kind_arch_isa != NULL)
2221 58 : for (tree p = properties; p && result; p = TREE_CHAIN (p))
2222 : {
2223 29 : const char *prop = omp_context_name_list_prop (p);
2224 29 : if (prop == NULL)
2225 : result = 0;
2226 : else
2227 29 : result = targetm.omp.device_kind_arch_isa (omp_device_arch,
2228 : prop);
2229 29 : if (dump_file)
2230 0 : fprintf (dump_file, "Matching device arch %s = %s\n",
2231 : prop, (result ? "true" : "false"));
2232 : }
2233 : else
2234 : {
2235 0 : result = 0;
2236 0 : if (dump_file)
2237 0 : fprintf (dump_file, "Cannot match device arch on target\n");
2238 : }
2239 : break;
2240 16 : case OMP_TRAIT_DEVICE_ISA:
2241 16 : if (targetm.omp.device_kind_arch_isa != NULL)
2242 32 : for (tree p = properties; p && result; p = TREE_CHAIN (p))
2243 : {
2244 16 : const char *prop = omp_context_name_list_prop (p);
2245 16 : if (prop == NULL)
2246 : result = 0;
2247 : else
2248 16 : result = targetm.omp.device_kind_arch_isa (omp_device_isa,
2249 : prop);
2250 16 : if (dump_file)
2251 0 : fprintf (dump_file, "Matching device isa %s = %s\n",
2252 : prop, (result ? "true" : "false"));
2253 : }
2254 : else
2255 : {
2256 0 : result = 0;
2257 0 : if (dump_file)
2258 0 : fprintf (dump_file, "Cannot match device isa on target\n");
2259 : }
2260 : break;
2261 0 : default:
2262 0 : gcc_unreachable ();
2263 : }
2264 153 : return result;
2265 : }
2266 :
2267 : /* Called for late resolution of the OMP_TARGET_DEVICE_MATCHES tree node to
2268 : a constant in omp-offload.cc. This is used in code that is wrapped in a
2269 : #pragma omp target construct to execute on the specified device, and
2270 : can be reduced to a compile-time constant in the offload compiler.
2271 : NODE is an OMP_TARGET_DEVICE_MATCHES tree node and the result is an
2272 : INTEGER_CST. */
2273 : tree
2274 0 : resolve_omp_target_device_matches (tree node)
2275 : {
2276 0 : tree sel = OMP_TARGET_DEVICE_MATCHES_SELECTOR (node);
2277 0 : enum omp_ts_code selector = (enum omp_ts_code) tree_to_shwi (sel);
2278 0 : tree properties = OMP_TARGET_DEVICE_MATCHES_PROPERTIES (node);
2279 0 : if (omp_target_device_matches_on_host (selector, properties))
2280 0 : return integer_one_node;
2281 : else
2282 0 : return integer_zero_node;
2283 : }
2284 :
2285 : /* Compare construct={simd} CLAUSES1 with CLAUSES2, return 0/-1/1/2 as
2286 : in omp_context_selector_set_compare. If MATCH_P is true, additionally
2287 : apply the special matching rules for the "simdlen" and "aligned" clauses
2288 : used to determine whether the selector CLAUSES1 is part of matches
2289 : the OpenMP context containing CLAUSES2. */
2290 :
2291 : static int
2292 58 : omp_construct_simd_compare (tree clauses1, tree clauses2, bool match_p)
2293 : {
2294 58 : if (clauses1 == NULL_TREE)
2295 0 : return clauses2 == NULL_TREE ? 0 : -1;
2296 58 : if (clauses2 == NULL_TREE)
2297 : return 1;
2298 :
2299 102 : int r = 0;
2300 68 : struct declare_variant_simd_data {
2301 : bool inbranch, notinbranch;
2302 : tree simdlen;
2303 : auto_vec<tree,16> data_sharing;
2304 : auto_vec<tree,16> aligned;
2305 68 : declare_variant_simd_data ()
2306 68 : : inbranch(false), notinbranch(false), simdlen(NULL_TREE) {}
2307 238 : } data[2];
2308 : unsigned int i;
2309 : tree e0, e1;
2310 102 : for (i = 0; i < 2; i++)
2311 332 : for (tree c = i ? clauses2 : clauses1; c; c = OMP_CLAUSE_CHAIN (c))
2312 : {
2313 196 : vec<tree> *v;
2314 196 : switch (OMP_CLAUSE_CODE (c))
2315 : {
2316 12 : case OMP_CLAUSE_INBRANCH:
2317 12 : data[i].inbranch = true;
2318 12 : continue;
2319 40 : case OMP_CLAUSE_NOTINBRANCH:
2320 40 : data[i].notinbranch = true;
2321 40 : continue;
2322 68 : case OMP_CLAUSE_SIMDLEN:
2323 68 : data[i].simdlen = OMP_CLAUSE_SIMDLEN_EXPR (c);
2324 68 : continue;
2325 52 : case OMP_CLAUSE_UNIFORM:
2326 52 : case OMP_CLAUSE_LINEAR:
2327 52 : v = &data[i].data_sharing;
2328 52 : break;
2329 24 : case OMP_CLAUSE_ALIGNED:
2330 24 : v = &data[i].aligned;
2331 24 : break;
2332 0 : default:
2333 0 : gcc_unreachable ();
2334 : }
2335 76 : unsigned HOST_WIDE_INT argno = tree_to_uhwi (OMP_CLAUSE_DECL (c));
2336 76 : if (argno >= v->length ())
2337 76 : v->safe_grow_cleared (argno + 1, true);
2338 76 : (*v)[argno] = c;
2339 : }
2340 : /* Here, r is used as a bitmask, 2 is set if CLAUSES1 has something
2341 : CLAUSES2 doesn't, 1 is set if CLAUSES2 has something CLAUSES1
2342 : doesn't. Thus, r == 3 implies return value 2, r == 1 implies
2343 : -1, r == 2 implies 1 and r == 0 implies 0. */
2344 34 : if (data[0].inbranch != data[1].inbranch)
2345 0 : r |= data[0].inbranch ? 2 : 1;
2346 34 : if (data[0].notinbranch != data[1].notinbranch)
2347 16 : r |= data[0].notinbranch ? 2 : 1;
2348 34 : e0 = data[0].simdlen;
2349 34 : e1 = data[1].simdlen;
2350 34 : if (!simple_cst_equal (e0, e1))
2351 : {
2352 8 : if (e0 && e1)
2353 : {
2354 8 : if (match_p && tree_fits_uhwi_p (e0) && tree_fits_uhwi_p (e1))
2355 : {
2356 : /* The two simdlen clauses match if m is a multiple of n. */
2357 8 : unsigned HOST_WIDE_INT n = tree_to_uhwi (e0);
2358 8 : unsigned HOST_WIDE_INT m = tree_to_uhwi (e1);
2359 8 : if (m % n != 0)
2360 : return 2;
2361 : }
2362 : else
2363 : return 2;
2364 : }
2365 0 : r |= data[0].simdlen ? 2 : 1;
2366 : }
2367 78 : if (data[0].data_sharing.length () < data[1].data_sharing.length ()
2368 52 : || data[0].aligned.length () < data[1].aligned.length ())
2369 0 : r |= 1;
2370 : tree c1, c2;
2371 90 : FOR_EACH_VEC_ELT (data[0].data_sharing, i, c1)
2372 : {
2373 72 : c2 = (i < data[1].data_sharing.length ()
2374 120 : ? data[1].data_sharing[i] : NULL_TREE);
2375 72 : if ((c1 == NULL_TREE) != (c2 == NULL_TREE))
2376 : {
2377 8 : r |= c1 != NULL_TREE ? 2 : 1;
2378 8 : continue;
2379 : }
2380 64 : if (c1 == NULL_TREE)
2381 46 : continue;
2382 18 : if (OMP_CLAUSE_CODE (c1) != OMP_CLAUSE_CODE (c2))
2383 : return 2;
2384 13 : if (OMP_CLAUSE_CODE (c1) != OMP_CLAUSE_LINEAR)
2385 7 : continue;
2386 6 : if (OMP_CLAUSE_LINEAR_VARIABLE_STRIDE (c1)
2387 6 : != OMP_CLAUSE_LINEAR_VARIABLE_STRIDE (c2))
2388 : return 2;
2389 6 : if (OMP_CLAUSE_LINEAR_KIND (c1) != OMP_CLAUSE_LINEAR_KIND (c2))
2390 : return 2;
2391 6 : if (!simple_cst_equal (OMP_CLAUSE_LINEAR_STEP (c1),
2392 6 : OMP_CLAUSE_LINEAR_STEP (c2)))
2393 : return 2;
2394 : }
2395 55 : FOR_EACH_VEC_ELT (data[0].aligned, i, c1)
2396 : {
2397 60 : c2 = i < data[1].aligned.length () ? data[1].aligned[i] : NULL_TREE;
2398 42 : if ((c1 == NULL_TREE) != (c2 == NULL_TREE))
2399 : {
2400 8 : r |= c1 != NULL_TREE ? 2 : 1;
2401 8 : continue;
2402 : }
2403 34 : if (c1 == NULL_TREE)
2404 28 : continue;
2405 6 : e0 = OMP_CLAUSE_ALIGNED_ALIGNMENT (c1);
2406 6 : e1 = OMP_CLAUSE_ALIGNED_ALIGNMENT (c2);
2407 6 : if (!simple_cst_equal (e0, e1))
2408 : {
2409 5 : if (e0 && e1
2410 1 : && match_p && tree_fits_uhwi_p (e0) && tree_fits_uhwi_p (e1))
2411 : {
2412 : /* The two aligned clauses match if n is a multiple of m. */
2413 0 : unsigned HOST_WIDE_INT n = tree_to_uhwi (e0);
2414 0 : unsigned HOST_WIDE_INT m = tree_to_uhwi (e1);
2415 0 : if (n % m != 0)
2416 : return 2;
2417 : }
2418 : else
2419 : return 2;
2420 : }
2421 : }
2422 13 : switch (r)
2423 : {
2424 : case 0: return 0;
2425 0 : case 1: return -1;
2426 8 : case 2: return 1;
2427 : case 3: return 2;
2428 0 : default: gcc_unreachable ();
2429 : }
2430 102 : }
2431 :
2432 : /* Compare properties of selectors SEL from SET other than construct.
2433 : CTX1 and CTX2 are the lists of properties to compare.
2434 : Return 0/-1/1/2 as in omp_context_selector_set_compare.
2435 : Unlike set names or selector names, properties can have duplicates. */
2436 :
2437 : static int
2438 171 : omp_context_selector_props_compare (enum omp_tss_code set,
2439 : enum omp_ts_code sel,
2440 : tree ctx1, tree ctx2)
2441 : {
2442 171 : int ret = 0;
2443 374 : for (int pass = 0; pass < 2; pass++)
2444 758 : for (tree p1 = pass ? ctx2 : ctx1; p1; p1 = TREE_CHAIN (p1))
2445 : {
2446 285 : tree p2;
2447 640 : for (p2 = pass ? ctx1 : ctx2; p2; p2 = TREE_CHAIN (p2))
2448 345 : if (OMP_TP_NAME (p1) == OMP_TP_NAME (p2))
2449 : {
2450 345 : if (OMP_TP_NAME (p1) == NULL_TREE)
2451 : {
2452 173 : if (set == OMP_TRAIT_SET_USER
2453 173 : && sel == OMP_TRAIT_USER_CONDITION)
2454 : {
2455 : /* Recognize constants that have equal truth values,
2456 : otherwise assume all expressions are unique. */
2457 122 : tree v1 = OMP_TP_VALUE (p1);
2458 122 : tree v2 = OMP_TP_VALUE (p2);
2459 122 : if (TREE_CODE (v1) != INTEGER_CST
2460 106 : || TREE_CODE (v2) != INTEGER_CST
2461 228 : || integer_zerop (v1) != integer_zerop (v2))
2462 16 : return 2;
2463 : break;
2464 : }
2465 51 : if (set == OMP_TRAIT_SET_TARGET_DEVICE
2466 51 : && sel == OMP_TRAIT_DEVICE_NUM)
2467 : {
2468 : /* Recognize constants that have equal values,
2469 : otherwise assume all expressions are unique. */
2470 51 : tree v1 = OMP_TP_VALUE (p1);
2471 51 : tree v2 = OMP_TP_VALUE (p2);
2472 51 : if (TREE_CODE (v1) != INTEGER_CST
2473 0 : || TREE_CODE (v2) != INTEGER_CST
2474 51 : || tree_int_cst_compare (v1, v2) != 0)
2475 51 : return 2;
2476 : break;
2477 : }
2478 0 : if (simple_cst_equal (OMP_TP_VALUE (p1), OMP_TP_VALUE (p2)))
2479 : break;
2480 : }
2481 172 : else if (OMP_TP_NAME (p1) == OMP_TP_NAMELIST_NODE)
2482 : {
2483 : /* Handle string constant vs identifier comparison for
2484 : name-list properties. */
2485 172 : const char *n1 = omp_context_name_list_prop (p1);
2486 172 : const char *n2 = omp_context_name_list_prop (p2);
2487 172 : if (n1 && n2 && !strcmp (n1, n2))
2488 : break;
2489 : }
2490 : else
2491 : break;
2492 : }
2493 208 : if (p2 == NULL_TREE)
2494 : {
2495 10 : int r = pass ? -1 : 1;
2496 10 : if (ret && ret != r)
2497 : return 2;
2498 10 : else if (pass)
2499 : return r;
2500 : else
2501 : {
2502 : ret = r;
2503 : break;
2504 : }
2505 : }
2506 : }
2507 : return ret;
2508 : }
2509 :
2510 : /* Compare single context selector sets CTX1 and CTX2 with SET name.
2511 : CTX1 and CTX2 are lists of trait-selectors.
2512 : Return 0 if CTX1 is equal to CTX2,
2513 : -1 if CTX1 is a strict subset of CTX2,
2514 : 1 if CTX2 is a strict subset of CTX1, or
2515 : 2 if neither context is a subset of another one. */
2516 :
2517 : static int
2518 2665 : omp_context_selector_set_compare (enum omp_tss_code set, tree ctx1, tree ctx2)
2519 : {
2520 :
2521 : /* If either list includes an ignored selector trait, neither can
2522 : be a subset of the other. */
2523 5585 : for (tree ts = ctx1; ts; ts = TREE_CHAIN (ts))
2524 2920 : if (OMP_TS_CODE (ts) == OMP_TRAIT_INVALID)
2525 : return 2;
2526 5612 : for (tree ts = ctx2; ts; ts = TREE_CHAIN (ts))
2527 2947 : if (OMP_TS_CODE (ts) == OMP_TRAIT_INVALID)
2528 : return 2;
2529 :
2530 2665 : bool swapped = false;
2531 2665 : int ret = 0;
2532 2665 : int len1 = list_length (ctx1);
2533 2665 : int len2 = list_length (ctx2);
2534 2665 : int cnt = 0;
2535 2665 : if (len1 < len2)
2536 : {
2537 76 : swapped = true;
2538 76 : std::swap (ctx1, ctx2);
2539 76 : std::swap (len1, len2);
2540 : }
2541 :
2542 2665 : if (set == OMP_TRAIT_SET_CONSTRUCT)
2543 : {
2544 : tree ts1;
2545 : tree ts2 = ctx2;
2546 : /* Handle construct set specially. In this case the order
2547 : of the selector matters too. */
2548 420 : for (ts1 = ctx1; ts1; ts1 = TREE_CHAIN (ts1))
2549 411 : if (OMP_TS_CODE (ts1) == OMP_TS_CODE (ts2))
2550 : {
2551 328 : int r = 0;
2552 328 : if (OMP_TS_CODE (ts1) == OMP_TRAIT_CONSTRUCT_SIMD)
2553 18 : r = omp_construct_simd_compare (OMP_TS_PROPERTIES (ts1),
2554 18 : OMP_TS_PROPERTIES (ts2),
2555 : false);
2556 328 : if (r == 2 || (ret && r && (ret < 0) != (r < 0)))
2557 13 : return 2;
2558 315 : if (ret == 0)
2559 234 : ret = r;
2560 315 : ts2 = TREE_CHAIN (ts2);
2561 315 : if (ts2 == NULL_TREE)
2562 : {
2563 233 : ts1 = TREE_CHAIN (ts1);
2564 233 : break;
2565 : }
2566 : }
2567 83 : else if (ret < 0)
2568 : return 2;
2569 : else
2570 : ret = 1;
2571 242 : if (ts2 != NULL_TREE)
2572 : return 2;
2573 233 : if (ts1 != NULL_TREE)
2574 : {
2575 36 : if (ret < 0)
2576 : return 2;
2577 : ret = 1;
2578 : }
2579 197 : if (ret == 0)
2580 : return 0;
2581 89 : return swapped ? -ret : ret;
2582 : }
2583 2613 : for (tree ts1 = ctx1; ts1; ts1 = TREE_CHAIN (ts1))
2584 : {
2585 2430 : enum omp_ts_code sel = OMP_TS_CODE (ts1);
2586 2430 : tree ts2;
2587 2547 : for (ts2 = ctx2; ts2; ts2 = TREE_CHAIN (ts2))
2588 2440 : if (sel == OMP_TS_CODE (ts2))
2589 : {
2590 2323 : tree score1 = OMP_TS_SCORE (ts1);
2591 2323 : tree score2 = OMP_TS_SCORE (ts2);
2592 2144 : if ((score1 && score2 && !simple_cst_equal (score1, score2))
2593 179 : || (score1 && !score2)
2594 2495 : || (!score1 && score2))
2595 2160 : return 2;
2596 :
2597 326 : int r = omp_context_selector_props_compare (set, OMP_TS_CODE (ts1),
2598 163 : OMP_TS_PROPERTIES (ts1),
2599 163 : OMP_TS_PROPERTIES (ts2));
2600 163 : if (r == 2 || (ret && r && (ret < 0) != (r < 0)))
2601 : return 2;
2602 96 : if (ret == 0)
2603 91 : ret = r;
2604 96 : cnt++;
2605 96 : break;
2606 : }
2607 203 : if (ts2 == NULL_TREE)
2608 : {
2609 107 : if (ret == -1)
2610 : return 2;
2611 : ret = 1;
2612 : }
2613 : }
2614 183 : if (cnt < len2)
2615 : return 2;
2616 91 : if (ret == 0)
2617 : return 0;
2618 20 : return swapped ? -ret : ret;
2619 : }
2620 :
2621 : /* Compare whole context selector specification CTX1 and CTX2.
2622 : Return 0 if CTX1 is equal to CTX2,
2623 : -1 if CTX1 is a strict subset of CTX2,
2624 : 1 if CTX2 is a strict subset of CTX1, or
2625 : 2 if neither context is a subset of another one. */
2626 :
2627 : static int
2628 3726 : omp_context_selector_compare (tree ctx1, tree ctx2)
2629 : {
2630 3726 : bool swapped = false;
2631 3726 : int ret = 0;
2632 3726 : int len1 = list_length (ctx1);
2633 3726 : int len2 = list_length (ctx2);
2634 3726 : int cnt = 0;
2635 3726 : if (len1 < len2)
2636 : {
2637 89 : swapped = true;
2638 89 : std::swap (ctx1, ctx2);
2639 89 : std::swap (len1, len2);
2640 : }
2641 5362 : for (tree tss1 = ctx1; tss1; tss1 = TREE_CHAIN (tss1))
2642 : {
2643 3971 : enum omp_tss_code set = OMP_TSS_CODE (tss1);
2644 3971 : tree tss2;
2645 5457 : for (tss2 = ctx2; tss2; tss2 = TREE_CHAIN (tss2))
2646 3997 : if (set == OMP_TSS_CODE (tss2))
2647 : {
2648 2511 : int r
2649 : = omp_context_selector_set_compare
2650 5022 : (set, OMP_TSS_TRAIT_SELECTORS (tss1),
2651 2511 : OMP_TSS_TRAIT_SELECTORS (tss2));
2652 2511 : if (r == 2 || (ret && r && (ret < 0) != (r < 0)))
2653 : return 2;
2654 191 : if (ret == 0)
2655 138 : ret = r;
2656 191 : cnt++;
2657 191 : break;
2658 : }
2659 1651 : if (tss2 == NULL_TREE)
2660 : {
2661 1460 : if (ret == -1)
2662 : return 2;
2663 : ret = 1;
2664 : }
2665 : }
2666 1391 : if (cnt < len2)
2667 : return 2;
2668 155 : if (ret == 0)
2669 : return 0;
2670 147 : return swapped ? -ret : ret;
2671 : }
2672 :
2673 : /* From context selector CTX, return trait-selector with name SEL in
2674 : trait-selector-set with name SET if any, or NULL_TREE if not found. */
2675 : tree
2676 7750 : omp_get_context_selector (tree ctx, enum omp_tss_code set,
2677 : enum omp_ts_code sel)
2678 : {
2679 14133 : for (tree tss = ctx; tss; tss = TREE_CHAIN (tss))
2680 8766 : if (OMP_TSS_CODE (tss) == set)
2681 6066 : for (tree ts = OMP_TSS_TRAIT_SELECTORS (tss); ts; ts = TREE_CHAIN (ts))
2682 4559 : if (OMP_TS_CODE (ts) == sel)
2683 : return ts;
2684 : return NULL_TREE;
2685 : }
2686 :
2687 : /* Similar, but returns the whole trait-selector list for SET in CTX. */
2688 : tree
2689 5985 : omp_get_context_selector_list (tree ctx, enum omp_tss_code set)
2690 : {
2691 10759 : for (tree tss = ctx; tss; tss = TREE_CHAIN (tss))
2692 6448 : if (OMP_TSS_CODE (tss) == set)
2693 1674 : return OMP_TSS_TRAIT_SELECTORS (tss);
2694 : return NULL_TREE;
2695 : }
2696 :
2697 : /* Map string S onto a trait selector set code. */
2698 : enum omp_tss_code
2699 4539 : omp_lookup_tss_code (const char * s)
2700 : {
2701 11854 : for (int i = 0; i < OMP_TRAIT_SET_LAST; i++)
2702 11827 : if (strcmp (s, omp_tss_map[i]) == 0)
2703 : return (enum omp_tss_code) i;
2704 : return OMP_TRAIT_SET_INVALID;
2705 : }
2706 :
2707 : /* Map string S onto a trait selector code for set SET. */
2708 : enum omp_ts_code
2709 5288 : omp_lookup_ts_code (enum omp_tss_code set, const char *s)
2710 : {
2711 5288 : unsigned int mask = 1 << set;
2712 55855 : for (int i = 0; i < OMP_TRAIT_LAST; i++)
2713 55766 : if ((mask & omp_ts_map[i].tss_mask) != 0
2714 13126 : && strcmp (s, omp_ts_map[i].name) == 0)
2715 : return (enum omp_ts_code) i;
2716 : return OMP_TRAIT_INVALID;
2717 : }
2718 :
2719 :
2720 : /* Return true if the selector CTX is dynamic. */
2721 : static bool
2722 3631 : omp_selector_is_dynamic (tree ctx)
2723 : {
2724 3631 : tree user_sel = omp_get_context_selector (ctx, OMP_TRAIT_SET_USER,
2725 : OMP_TRAIT_USER_CONDITION);
2726 3631 : if (user_sel)
2727 : {
2728 1455 : tree expr = OMP_TP_VALUE (OMP_TS_PROPERTIES (user_sel));
2729 :
2730 : /* The user condition is not dynamic if it is constant. */
2731 1455 : if (!tree_fits_shwi_p (expr))
2732 : return true;
2733 : }
2734 :
2735 2470 : tree target_device_ss
2736 2470 : = omp_get_context_selector_list (ctx, OMP_TRAIT_SET_TARGET_DEVICE);
2737 2470 : if (target_device_ss)
2738 : return true;
2739 :
2740 : return false;
2741 : }
2742 :
2743 : /* Helper function for omp_dynamic_cond: return a boolean tree expression
2744 : that tests whether *DEVICE_NUM is a "conforming device number other
2745 : than omp_invalid_device". This may modify *DEVICE_NUM (i.e, to be
2746 : a save_expr). *IS_HOST is set to true if the device can be statically
2747 : determined to be the host. */
2748 :
2749 : static tree
2750 103 : omp_device_num_check (tree *device_num, bool *is_host)
2751 : {
2752 : /* C++ may wrap the device_num expr in a CLEANUP_POINT_EXPR; we want
2753 : to look inside of it for the special cases. */
2754 103 : tree t = *device_num;
2755 103 : if (TREE_CODE (t) == CLEANUP_POINT_EXPR)
2756 15 : t = TREE_OPERAND (t, 0);
2757 :
2758 : /* First check for some constant values we can treat specially. */
2759 103 : if (tree_fits_shwi_p (t))
2760 : {
2761 32 : HOST_WIDE_INT num = tree_to_shwi (t);
2762 32 : if (num < -1)
2763 4 : return integer_zero_node;
2764 : /* Initial device? */
2765 28 : if (num == -1)
2766 : {
2767 20 : *is_host = true;
2768 20 : return integer_one_node;
2769 : }
2770 : /* There is always at least one device (the host + offload devices). */
2771 8 : if (num == 0)
2772 4 : return integer_one_node;
2773 : /* If there is no offloading, there is exactly one device. */
2774 4 : if (!ENABLE_OFFLOADING && num > 0)
2775 4 : return integer_zero_node;
2776 : }
2777 :
2778 : /* Also test for direct calls to OpenMP routines that return valid
2779 : device numbers. */
2780 71 : if (TREE_CODE (t) == CALL_EXPR)
2781 : {
2782 20 : tree fndecl = get_callee_fndecl (t);
2783 20 : if (fndecl && omp_runtime_api_call (fndecl))
2784 : {
2785 20 : const char *fnname = IDENTIFIER_POINTER (DECL_NAME (fndecl));
2786 20 : if (strcmp (fnname, "omp_get_default_device") == 0
2787 12 : || strcmp (fnname, "omp_get_device_num") == 0)
2788 12 : return integer_one_node;
2789 8 : if (strcmp (fnname, "omp_get_num_devices") == 0
2790 4 : || strcmp (fnname, "omp_get_initial_device") == 0)
2791 : {
2792 8 : *is_host = true;
2793 8 : return integer_one_node;
2794 : }
2795 : }
2796 : }
2797 :
2798 : /* Otherwise, test that -1 <= *device_num <= omp_get_num_devices (). */
2799 51 : *device_num = save_expr (*device_num);
2800 51 : tree lotest = build2 (GE_EXPR, integer_type_node, *device_num,
2801 : integer_minus_one_node);
2802 51 : tree fndecl = builtin_decl_explicit (BUILT_IN_OMP_GET_NUM_DEVICES);
2803 51 : tree hitest = build2 (LE_EXPR, integer_type_node, *device_num,
2804 : build_call_expr (fndecl, 0));
2805 51 : return build2 (TRUTH_ANDIF_EXPR, integer_type_node, lotest, hitest);
2806 : }
2807 :
2808 : /* Return a tree expression representing the dynamic part of the context
2809 : selector CTX. SUPERCONTEXT is the surrounding BLOCK, in case we need
2810 : to introduce a new BLOCK in the result. */
2811 : tree
2812 513 : omp_dynamic_cond (tree ctx, tree supercontext)
2813 : {
2814 513 : tree user_cond = NULL_TREE, target_device_cond = NULL_TREE;
2815 :
2816 : /* Build the "user" part of the dynamic selector. This is a test
2817 : predicate taken directly for the "condition" trait in this set. */
2818 513 : tree user_sel = omp_get_context_selector (ctx, OMP_TRAIT_SET_USER,
2819 : OMP_TRAIT_USER_CONDITION);
2820 513 : if (user_sel)
2821 : {
2822 163 : tree expr = OMP_TP_VALUE (OMP_TS_PROPERTIES (user_sel));
2823 :
2824 : /* The user condition is not dynamic if it is constant. */
2825 163 : if (!tree_fits_shwi_p (expr))
2826 161 : user_cond = unshare_expr (expr);
2827 : }
2828 :
2829 : /* Build the "target_device" part of the dynamic selector. In the
2830 : most general case this requires building a bit of code that runs
2831 : on the specified device_num using the same mechanism as
2832 : "#pragma omp target" that uses the OMP_TARGET_DEVICE_MATCHES magic
2833 : cookie to represent the kind/arch/isa tests which are and'ed together.
2834 : These cookies can be resolved into a constant truth value by the
2835 : offload compiler; see resolve_omp_target_device_matches, above.
2836 :
2837 : In some cases, we can (in)validate the device number in advance.
2838 : If it is not valid, the whole selector fails to match. If it is
2839 : valid and refers to the host (e.g., constant -1), then we can
2840 : resolve the match to a constant truth value now instead of having
2841 : to create a OMP_TARGET_DEVICE_MATCHES. */
2842 :
2843 513 : tree target_device_ss
2844 513 : = omp_get_context_selector_list (ctx, OMP_TRAIT_SET_TARGET_DEVICE);
2845 513 : if (target_device_ss)
2846 : {
2847 138 : tree device_num = NULL_TREE;
2848 138 : tree kind = NULL_TREE;
2849 138 : tree arch = NULL_TREE;
2850 138 : tree isa = NULL_TREE;
2851 138 : tree device_ok = NULL_TREE;
2852 138 : bool is_host = !ENABLE_OFFLOADING;
2853 :
2854 138 : tree device_num_sel
2855 138 : = omp_get_context_selector (ctx, OMP_TRAIT_SET_TARGET_DEVICE,
2856 : OMP_TRAIT_DEVICE_NUM);
2857 138 : if (device_num_sel)
2858 : {
2859 103 : device_num = OMP_TP_VALUE (OMP_TS_PROPERTIES (device_num_sel));
2860 103 : device_ok = omp_device_num_check (&device_num, &is_host);
2861 : /* If an invalid constant device number was specified, the
2862 : whole selector fails to match, and there's no point in
2863 : continuing to generate code that would never be executed. */
2864 103 : if (device_ok == integer_zero_node)
2865 : {
2866 8 : target_device_cond = integer_zero_node;
2867 51 : goto wrapup;
2868 : }
2869 : }
2870 :
2871 130 : tree kind_sel
2872 130 : = omp_get_context_selector (ctx, OMP_TRAIT_SET_TARGET_DEVICE,
2873 : OMP_TRAIT_DEVICE_KIND);
2874 : /* "any" is equivalent to omitting this trait selector. */
2875 130 : if (kind_sel
2876 240 : && strcmp (omp_context_name_list_prop (OMP_TS_PROPERTIES (kind_sel)),
2877 : "any"))
2878 : {
2879 108 : tree props = OMP_TS_PROPERTIES (kind_sel);
2880 108 : if (!is_host)
2881 0 : kind = build2 (OMP_TARGET_DEVICE_MATCHES, integer_type_node,
2882 : build_int_cst (integer_type_node,
2883 : (int) OMP_TRAIT_DEVICE_KIND),
2884 : props);
2885 108 : else if (!omp_target_device_matches_on_host (OMP_TRAIT_DEVICE_KIND,
2886 : props))
2887 : {
2888 : /* The whole selector fails to match. */
2889 43 : target_device_cond = integer_zero_node;
2890 43 : goto wrapup;
2891 : }
2892 : /* else it is statically resolved to true and is a no-op. */
2893 : }
2894 87 : tree arch_sel
2895 87 : = omp_get_context_selector (ctx, OMP_TRAIT_SET_TARGET_DEVICE,
2896 : OMP_TRAIT_DEVICE_ARCH);
2897 87 : if (arch_sel)
2898 : {
2899 29 : tree props = OMP_TS_PROPERTIES (arch_sel);
2900 29 : if (!is_host)
2901 0 : arch = build2 (OMP_TARGET_DEVICE_MATCHES, integer_type_node,
2902 : build_int_cst (integer_type_node,
2903 : (int) OMP_TRAIT_DEVICE_ARCH),
2904 : props);
2905 29 : else if (!omp_target_device_matches_on_host (OMP_TRAIT_DEVICE_ARCH,
2906 : props))
2907 : {
2908 : /* The whole selector fails to match. */
2909 0 : target_device_cond = integer_zero_node;
2910 0 : goto wrapup;
2911 : }
2912 : /* else it is statically resolved to true and is a no-op. */
2913 : }
2914 :
2915 87 : tree isa_sel
2916 87 : = omp_get_context_selector (ctx, OMP_TRAIT_SET_TARGET_DEVICE,
2917 : OMP_TRAIT_DEVICE_ISA);
2918 87 : if (isa_sel)
2919 : {
2920 16 : tree props = OMP_TS_PROPERTIES (isa_sel);
2921 16 : if (!is_host)
2922 0 : isa = build2 (OMP_TARGET_DEVICE_MATCHES, integer_type_node,
2923 : build_int_cst (integer_type_node,
2924 : (int) OMP_TRAIT_DEVICE_ISA),
2925 : props);
2926 16 : else if (!omp_target_device_matches_on_host (OMP_TRAIT_DEVICE_ISA,
2927 : props))
2928 : {
2929 : /* The whole selector fails to match. */
2930 0 : target_device_cond = integer_zero_node;
2931 0 : goto wrapup;
2932 : }
2933 : /* else it is statically resolved to true and is a no-op. */
2934 : }
2935 :
2936 : /* AND the three possible tests together. */
2937 87 : tree test_expr = kind ? kind : NULL_TREE;
2938 87 : if (arch && test_expr)
2939 0 : test_expr = build2 (TRUTH_ANDIF_EXPR, integer_type_node,
2940 : arch, test_expr);
2941 87 : else if (arch)
2942 0 : test_expr = arch;
2943 87 : if (isa && test_expr)
2944 0 : test_expr = build2 (TRUTH_ANDIF_EXPR, integer_type_node,
2945 : isa, test_expr);
2946 87 : else if (isa)
2947 : test_expr = isa;
2948 :
2949 87 : if (!test_expr)
2950 : /* This could happen if the selector includes only kind="any",
2951 : or is_host is true and it could be statically determined to
2952 : be true. The selector always matches, but we still have to
2953 : evaluate the device_num expression. */
2954 : {
2955 87 : if (device_num)
2956 67 : target_device_cond = build2 (COMPOUND_EXPR, integer_type_node,
2957 : device_num, integer_one_node);
2958 : else
2959 20 : target_device_cond = integer_one_node;
2960 : }
2961 : else
2962 : {
2963 : /* Arrange to evaluate test_expr in the offload compiler for
2964 : device device_num. */
2965 0 : tree stmt = make_node (OMP_TARGET);
2966 0 : TREE_TYPE (stmt) = void_type_node;
2967 0 : tree result_var = create_tmp_var (integer_type_node, "td_match");
2968 0 : tree map = build_omp_clause (UNKNOWN_LOCATION, OMP_CLAUSE_MAP);
2969 0 : OMP_CLAUSE_DECL (map) = result_var;
2970 0 : OMP_CLAUSE_SET_MAP_KIND (map, GOMP_MAP_FROM);
2971 0 : OMP_TARGET_CLAUSES (stmt) = map;
2972 0 : if (device_num)
2973 : {
2974 0 : tree clause = build_omp_clause (UNKNOWN_LOCATION,
2975 : OMP_CLAUSE_DEVICE);
2976 0 : OMP_CLAUSE_CHAIN (clause) = NULL_TREE;
2977 0 : OMP_CLAUSE_DEVICE_ID (clause) = device_num;
2978 0 : OMP_CLAUSE_DEVICE_ANCESTOR (clause) = false;
2979 0 : OMP_CLAUSE_CHAIN (map) = clause;
2980 : }
2981 :
2982 0 : tree block = make_node (BLOCK);
2983 0 : BLOCK_SUPERCONTEXT (block) = supercontext;
2984 :
2985 0 : tree bind = build3 (BIND_EXPR, void_type_node, NULL_TREE,
2986 : build2 (MODIFY_EXPR, integer_type_node,
2987 : result_var, test_expr),
2988 : block);
2989 0 : TREE_SIDE_EFFECTS (bind) = 1;
2990 0 : OMP_TARGET_BODY (stmt) = bind;
2991 0 : target_device_cond = build2 (COMPOUND_EXPR, integer_type_node,
2992 : stmt, result_var);
2993 :
2994 : /* If necessary, "and" target_device_cond with the test to
2995 : make sure the device number is valid. */
2996 0 : if (device_ok && device_ok != integer_one_node)
2997 0 : target_device_cond = build2 (TRUTH_ANDIF_EXPR, integer_type_node,
2998 : device_ok, target_device_cond);
2999 :
3000 : /* Set the bit to trigger resolution of OMP_TARGET_DEVICE_MATCHES
3001 : in the ompdevlow pass. */
3002 0 : if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
3003 0 : cgraph_node::get (cfun->decl)->has_omp_variant_constructs = 1;
3004 : }
3005 : }
3006 :
3007 375 : wrapup:
3008 513 : if (user_cond && target_device_cond)
3009 12 : return build2 (TRUTH_ANDIF_EXPR, integer_type_node,
3010 12 : user_cond, target_device_cond);
3011 501 : else if (user_cond)
3012 149 : return user_cond;
3013 : else if (target_device_cond)
3014 : return target_device_cond;
3015 : else
3016 : return NULL_TREE;
3017 : }
3018 :
3019 :
3020 : /* Given an omp_variant VARIANT, compute VARIANT->score and
3021 : VARIANT->scorable.
3022 : CONSTRUCT_CONTEXT is the OpenMP construct context; if this is null or
3023 : COMPLETE_P is false (e.g., during parsing or gimplification) then it
3024 : may not be possible to compute the score accurately and the scorable
3025 : flag is set to false.
3026 :
3027 : Cited text in the comments is from section 7.2 of the OpenMP 5.2
3028 : specification. */
3029 :
3030 : static void
3031 3631 : omp_context_compute_score (struct omp_variant *variant,
3032 : tree construct_context, bool complete_p)
3033 : {
3034 3631 : int l = list_length (construct_context);
3035 3631 : tree ctx = variant->selector;
3036 3631 : variant->scorable = true;
3037 :
3038 : /* "the final score is the sum of the values of all specified selectors
3039 : plus 1". */
3040 3631 : variant->score = 1;
3041 7709 : for (tree tss = ctx; tss; tss = TREE_CHAIN (tss))
3042 : {
3043 4078 : if (OMP_TSS_CODE (tss) == OMP_TRAIT_SET_CONSTRUCT)
3044 : {
3045 : /* "Each trait selector for which the corresponding trait appears
3046 : in the context trait set in the OpenMP context..." */
3047 1462 : score_wide_int tss_score = 0;
3048 1462 : omp_construct_traits_match (OMP_TSS_TRAIT_SELECTORS (tss),
3049 : construct_context, &tss_score);
3050 1462 : variant->score += tss_score;
3051 1462 : if (!complete_p)
3052 4 : variant->scorable = false;
3053 : }
3054 2616 : else if (OMP_TSS_CODE (tss) == OMP_TRAIT_SET_DEVICE
3055 2616 : || OMP_TSS_CODE (tss) == OMP_TRAIT_SET_TARGET_DEVICE)
3056 : {
3057 : /* "The kind, arch, and isa selectors, if specified, are given
3058 : the values 2**l, 2**(l+1), and 2**(l+2), respectively..."
3059 : FIXME: the spec isn't clear what should happen if there are
3060 : both "device" and "target_device" selector sets specified.
3061 : This implementation adds up the bits rather than ORs them. */
3062 1460 : for (tree ts = OMP_TSS_TRAIT_SELECTORS (tss); ts;
3063 858 : ts = TREE_CHAIN (ts))
3064 : {
3065 858 : enum omp_ts_code code = OMP_TS_CODE (ts);
3066 858 : if (code == OMP_TRAIT_DEVICE_KIND)
3067 461 : variant->score
3068 461 : += wi::shifted_mask <score_wide_int> (l, 1, false);
3069 397 : else if (code == OMP_TRAIT_DEVICE_ARCH)
3070 119 : variant->score
3071 119 : += wi::shifted_mask <score_wide_int> (l + 1, 1, false);
3072 278 : else if (code == OMP_TRAIT_DEVICE_ISA)
3073 151 : variant->score
3074 151 : += wi::shifted_mask <score_wide_int> (l + 2, 1, false);
3075 : }
3076 602 : if (!complete_p)
3077 195 : variant->scorable = false;
3078 : }
3079 : else
3080 : {
3081 : /* "Trait selectors for which a trait-score is specified..."
3082 : Note that there are no implementation-defined selectors, and
3083 : "other selectors are given a value of zero". */
3084 4028 : for (tree ts = OMP_TSS_TRAIT_SELECTORS (tss); ts;
3085 2014 : ts = TREE_CHAIN (ts))
3086 : {
3087 2014 : tree s = OMP_TS_SCORE (ts);
3088 1436 : if (s && TREE_CODE (s) == INTEGER_CST)
3089 1436 : variant->score
3090 2872 : += score_wide_int::from (wi::to_wide (s),
3091 2872 : TYPE_SIGN (TREE_TYPE (s)));
3092 : }
3093 : }
3094 : }
3095 3631 : }
3096 :
3097 : /* CONSTRUCT_CONTEXT contains "the directive names, each being a trait,
3098 : of all enclosing constructs at that point in the program up to a target
3099 : construct", per section 7.1 of the 5.2 specification. The traits are
3100 : collected during gimplification and are listed outermost first.
3101 :
3102 : This function attempts to apply the "if the point in the program is not
3103 : enclosed by a target construct, the following rules are applied in order"
3104 : requirements that follow in the same paragraph. This may not be possible,
3105 : depending on the compilation phase; in particular, "declare simd" clones
3106 : are not known until late resolution.
3107 :
3108 : The augmented context is returned, and *COMPLETEP is set to true if
3109 : the context is known to be complete, false otherwise. */
3110 : static tree
3111 5370 : omp_complete_construct_context (tree construct_context, bool *completep)
3112 : {
3113 : /* The point in the program is enclosed by a target construct. */
3114 5370 : if (construct_context
3115 8566 : && OMP_TS_CODE (construct_context) == OMP_TRAIT_CONSTRUCT_TARGET)
3116 242 : *completep = true;
3117 :
3118 : /* At parse time we have none of the information we need to collect
3119 : the missing pieces. */
3120 5128 : else if (symtab->state == PARSING)
3121 640 : *completep = false;
3122 :
3123 : else
3124 : {
3125 4488 : tree attributes = DECL_ATTRIBUTES (current_function_decl);
3126 :
3127 : /* Add simd trait when in a simd clone. This information is only
3128 : available during late resolution in the omp_device_lower pass,
3129 : however we can also rule out cases where we know earlier that
3130 : cfun is not a candidate for cloning. */
3131 4488 : if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
3132 : {
3133 304 : cgraph_node *node = cgraph_node::get (cfun->decl);
3134 304 : if (node->simdclone)
3135 288 : construct_context = make_trait_selector (OMP_TRAIT_CONSTRUCT_SIMD,
3136 : NULL_TREE, NULL_TREE,
3137 : construct_context);
3138 304 : *completep = true;
3139 304 : }
3140 4184 : else if (lookup_attribute ("omp declare simd", attributes))
3141 42 : *completep = false;
3142 : else
3143 4142 : *completep = true;
3144 :
3145 : /* Add construct selector set within a "declare variant" function. */
3146 4488 : tree variant_attr
3147 4488 : = lookup_attribute ("omp declare variant variant", attributes);
3148 4488 : if (variant_attr)
3149 : {
3150 28 : tree temp = NULL_TREE;
3151 48 : for (tree t = TREE_VALUE (variant_attr); t; t = TREE_CHAIN (t))
3152 20 : temp = chainon (temp, copy_node (t));
3153 28 : construct_context = chainon (temp, construct_context);
3154 : }
3155 :
3156 : /* Add target trait when in a target variant. */
3157 4488 : if (lookup_attribute ("omp declare target", attributes))
3158 74 : construct_context = make_trait_selector (OMP_TRAIT_CONSTRUCT_TARGET,
3159 : NULL_TREE, NULL_TREE,
3160 : construct_context);
3161 : }
3162 5370 : return construct_context;
3163 : }
3164 :
3165 : /* Comparison function for sorting routines, to sort OpenMP metadirective
3166 : variants by decreasing score. */
3167 :
3168 : static int
3169 11991 : sort_variant (const void * a, const void *b, void *)
3170 : {
3171 11991 : score_wide_int score1
3172 : = ((const struct omp_variant *) a)->score;
3173 11991 : score_wide_int score2
3174 : = ((const struct omp_variant *) b)->score;
3175 :
3176 11991 : if (score1 > score2)
3177 : return -1;
3178 5240 : else if (score1 < score2)
3179 : return 1;
3180 : else
3181 279 : return 0;
3182 : }
3183 :
3184 : /* Return a vector of dynamic replacement candidates for the directive
3185 : candidates in ALL_VARIANTS. Return an empty vector if the candidates
3186 : cannot be resolved. */
3187 :
3188 : vec<struct omp_variant>
3189 2837 : omp_get_dynamic_candidates (vec <struct omp_variant> &all_variants,
3190 : tree construct_context)
3191 : {
3192 2837 : auto_vec <struct omp_variant> variants;
3193 2837 : struct omp_variant default_variant;
3194 2837 : bool default_found = false;
3195 2837 : bool complete_p;
3196 :
3197 2837 : construct_context
3198 2837 : = omp_complete_construct_context (construct_context, &complete_p);
3199 :
3200 2837 : if (dump_file)
3201 : {
3202 78 : fprintf (dump_file, "\nIn omp_get_dynamic_candidates:\n");
3203 78 : if (symtab->state == PARSING)
3204 78 : fprintf (dump_file, "invoked during parsing\n");
3205 0 : else if (cfun && (cfun->curr_properties & PROP_gimple_any) == 0)
3206 0 : fprintf (dump_file, "invoked during gimplification\n");
3207 0 : else if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
3208 0 : fprintf (dump_file, "invoked during late resolution\n");
3209 : else
3210 0 : fprintf (dump_file, "confused about invocation context?!?\n");
3211 156 : fprintf (dump_file, "construct_context has %d traits (%s)\n",
3212 0 : (construct_context ? list_length (construct_context) : 0),
3213 : (complete_p ? "complete" : "incomplete"));
3214 : }
3215 :
3216 9315 : for (unsigned int i = 0; i < all_variants.length (); i++)
3217 : {
3218 6556 : struct omp_variant variant = all_variants[i];
3219 :
3220 6556 : if (variant.selector == NULL_TREE)
3221 : {
3222 2759 : gcc_assert (!default_found);
3223 2759 : default_found = true;
3224 2759 : default_variant = variant;
3225 2759 : default_variant.score = 0;
3226 2759 : default_variant.scorable = true;
3227 2759 : default_variant.matchable = true;
3228 2759 : default_variant.dynamic_selector = false;
3229 2759 : if (dump_file)
3230 54 : fprintf (dump_file,
3231 : "Considering default selector as candidate\n");
3232 2759 : continue;
3233 : }
3234 :
3235 3797 : variant.matchable = true;
3236 3797 : variant.scorable = true;
3237 :
3238 3797 : if (dump_file)
3239 : {
3240 78 : fprintf (dump_file, "Considering selector ");
3241 78 : print_omp_context_selector (dump_file, variant.selector, TDF_NONE);
3242 78 : fprintf (dump_file, " as candidate - ");
3243 : }
3244 :
3245 3797 : switch (omp_context_selector_matches (variant.selector,
3246 : construct_context, complete_p))
3247 : {
3248 99 : case -1:
3249 99 : if (dump_file)
3250 24 : fprintf (dump_file, "unmatchable\n");
3251 : /* At parse time, just give up if we can't determine whether
3252 : things match. */
3253 99 : if (symtab->state == PARSING)
3254 : {
3255 78 : variants.truncate (0);
3256 78 : return variants.copy ();
3257 : }
3258 : /* Otherwise we must be invoked from the gimplifier. */
3259 21 : gcc_assert (cfun && (cfun->curr_properties & PROP_gimple_any) == 0);
3260 21 : variant.matchable = false;
3261 : /* FALLTHRU */
3262 3631 : case 1:
3263 3631 : omp_context_compute_score (&variant, construct_context, complete_p);
3264 3631 : variant.dynamic_selector
3265 3631 : = omp_selector_is_dynamic (variant.selector);
3266 3631 : variants.safe_push (variant);
3267 3631 : if (dump_file && variant.matchable)
3268 : {
3269 54 : if (variant.dynamic_selector)
3270 48 : fprintf (dump_file, "matched, dynamic");
3271 : else
3272 6 : fprintf (dump_file, "matched, non-dynamic");
3273 : }
3274 : break;
3275 88 : case 0:
3276 88 : if (dump_file)
3277 0 : fprintf (dump_file, "no match");
3278 : break;
3279 : }
3280 :
3281 3719 : if (dump_file)
3282 54 : fprintf (dump_file, "\n");
3283 : }
3284 :
3285 : /* There must be one default variant. */
3286 2759 : gcc_assert (default_found);
3287 :
3288 : /* If there are no matching selectors, return the default. */
3289 2759 : if (variants.length () == 0)
3290 : {
3291 685 : variants.safe_push (default_variant);
3292 685 : return variants.copy ();
3293 : }
3294 :
3295 : /* If there is only one matching selector, use it. */
3296 2074 : if (variants.length () == 1)
3297 : {
3298 1460 : if (variants[0].matchable)
3299 : {
3300 1460 : if (variants[0].dynamic_selector)
3301 198 : variants.safe_push (default_variant);
3302 1460 : return variants.copy ();
3303 : }
3304 : else
3305 : {
3306 : /* We don't know whether the one non-default selector will
3307 : actually match. */
3308 0 : variants.truncate (0);
3309 0 : return variants.copy ();
3310 : }
3311 : }
3312 :
3313 : /* A context selector that is a strict subset of another context selector
3314 : has a score of zero. This only applies if the selector that is a
3315 : superset definitely matches, though. */
3316 2785 : for (unsigned int i = 0; i < variants.length (); i++)
3317 5811 : for (unsigned int j = i + 1; j < variants.length (); j++)
3318 : {
3319 7386 : int r = omp_context_selector_compare (variants[i].selector,
3320 3693 : variants[j].selector);
3321 3746 : if (r == -1 && variants[j].matchable)
3322 : {
3323 : /* variant i is a strict subset of variant j. */
3324 53 : variants[i].score = 0;
3325 53 : variants[i].scorable = true;
3326 53 : break;
3327 : }
3328 3734 : else if (r == 1 && variants[i].matchable)
3329 : /* variant j is a strict subset of variant i. */
3330 : {
3331 94 : variants[j].score = 0;
3332 94 : variants[j].scorable = true;
3333 : }
3334 : }
3335 :
3336 : /* Sort the variants by decreasing score, preserving the original order
3337 : in case of a tie. */
3338 614 : variants.stablesort (sort_variant, NULL);
3339 :
3340 : /* Add the default as a final choice. */
3341 614 : variants.safe_push (default_variant);
3342 :
3343 614 : if (dump_file)
3344 : {
3345 0 : fprintf (dump_file, "Sorted variants are:\n");
3346 0 : for (unsigned i = 0; i < variants.length (); i++)
3347 : {
3348 0 : HOST_WIDE_INT score = variants[i].score.to_shwi ();
3349 0 : fprintf (dump_file, "score %d matchable %d scorable %d ",
3350 0 : (int)score, (int)(variants[i].matchable),
3351 0 : (int)(variants[i].scorable));
3352 0 : if (variants[i].selector)
3353 : {
3354 0 : fprintf (dump_file, "selector ");
3355 0 : print_omp_context_selector (dump_file, variants[i].selector,
3356 : TDF_NONE);
3357 0 : fprintf (dump_file, "\n");
3358 : }
3359 : else
3360 0 : fprintf (dump_file, "default selector\n");
3361 : }
3362 : }
3363 :
3364 : /* Build the dynamic candidate list. */
3365 1711 : for (unsigned i = 0; i < variants.length (); i++)
3366 : {
3367 : /* If we encounter a candidate that wasn't definitely matched,
3368 : give up now. */
3369 1711 : if (!variants[i].matchable)
3370 : {
3371 4 : variants.truncate (0);
3372 4 : break;
3373 : }
3374 :
3375 : /* In general, we can't proceed if we can't accurately score any
3376 : of the selectors, since the sorting may be incorrect. But, since
3377 : the actual score will never be lower than the guessed value, we
3378 : can use the first variant if it is not scorable but either the next
3379 : one is a subset of the first, is scorable, or we can make a
3380 : direct comparison of the high-order isa/arch/kind bits. */
3381 1707 : if (!variants[i].scorable)
3382 : {
3383 28 : bool ok = true;
3384 28 : if (i != 0)
3385 : ok = false;
3386 20 : else if (variants[i+1].scorable)
3387 : /* ok */
3388 : ;
3389 20 : else if (variants[i+1].score > 0)
3390 : {
3391 : /* To keep comparisons simple, reject selectors that contain
3392 : sets other than device, target_device, or construct. */
3393 20 : for (tree tss = variants[i].selector;
3394 40 : tss && ok; tss = TREE_CHAIN (tss))
3395 : {
3396 20 : enum omp_tss_code code = OMP_TSS_CODE (tss);
3397 20 : if (code != OMP_TRAIT_SET_DEVICE
3398 : && code != OMP_TRAIT_SET_TARGET_DEVICE
3399 20 : && code != OMP_TRAIT_SET_CONSTRUCT)
3400 12 : ok = false;
3401 : }
3402 20 : for (tree tss = variants[i+1].selector;
3403 28 : tss && ok; tss = TREE_CHAIN (tss))
3404 : {
3405 8 : enum omp_tss_code code = OMP_TSS_CODE (tss);
3406 8 : if (code != OMP_TRAIT_SET_DEVICE
3407 : && code != OMP_TRAIT_SET_TARGET_DEVICE
3408 8 : && code != OMP_TRAIT_SET_CONSTRUCT)
3409 0 : ok = false;
3410 : }
3411 : /* Ignore the construct bits of the score. If the isa/arch/kind
3412 : bits are strictly ordered, we're good to go. Since
3413 : "the final score is the sum of the values of all specified
3414 : selectors plus 1", subtract that 1 from both scores before
3415 : getting rid of the low bits. */
3416 20 : if (ok)
3417 : {
3418 8 : size_t l = list_length (construct_context);
3419 8 : gcc_assert (variants[i].score > 0
3420 : && variants[i+1].score > 0);
3421 8 : if ((variants[i].score - 1) >> l
3422 16 : <= (variants[i+1].score - 1) >> l)
3423 0 : ok = false;
3424 : }
3425 : }
3426 :
3427 8 : if (!ok)
3428 : {
3429 20 : variants.truncate (0);
3430 20 : break;
3431 : }
3432 : }
3433 :
3434 1687 : if (dump_file)
3435 : {
3436 0 : fprintf (dump_file, "Adding directive variant with ");
3437 :
3438 0 : if (variants[i].selector)
3439 : {
3440 0 : fprintf (dump_file, "selector ");
3441 0 : print_omp_context_selector (dump_file, variants[i].selector,
3442 : TDF_NONE);
3443 : }
3444 : else
3445 0 : fprintf (dump_file, "default selector");
3446 :
3447 0 : fprintf (dump_file, " as candidate.\n");
3448 : }
3449 :
3450 : /* The last of the candidates is ended by a static selector. */
3451 1687 : if (!variants[i].dynamic_selector)
3452 : {
3453 590 : variants.truncate (i + 1);
3454 590 : break;
3455 : }
3456 : }
3457 :
3458 614 : return variants.copy ();
3459 2837 : }
3460 :
3461 : /* Two attempts are made to resolve calls to "declare variant" functions:
3462 : early resolution in the gimplifier, and late resolution in the
3463 : omp_device_lower pass. If early resolution is not possible, the
3464 : original function call is gimplified into the same form as metadirective
3465 : and goes through the same late resolution code as metadirective. */
3466 :
3467 : /* Collect "declare variant" candidates for BASE. CONSTRUCT_CONTEXT
3468 : is the un-augmented context, or NULL_TREE if that information is not
3469 : available yet. */
3470 : vec<struct omp_variant>
3471 2039 : omp_declare_variant_candidates (tree base, tree construct_context)
3472 : {
3473 2039 : auto_vec <struct omp_variant> candidates;
3474 2039 : bool complete_p;
3475 2039 : tree augmented_context
3476 2039 : = omp_complete_construct_context (construct_context, &complete_p);
3477 :
3478 : /* The variants are stored on (possible multiple) "omp declare variant base"
3479 : attributes on the base function. */
3480 4320 : for (tree attr = DECL_ATTRIBUTES (base); attr; attr = TREE_CHAIN (attr))
3481 : {
3482 2342 : attr = lookup_attribute ("omp declare variant base", attr);
3483 2342 : if (attr == NULL_TREE)
3484 : break;
3485 :
3486 2281 : tree fndecl = TREE_PURPOSE (TREE_VALUE (attr));
3487 2281 : tree selector = TREE_VALUE (TREE_VALUE (attr));
3488 :
3489 2281 : if (TREE_CODE (fndecl) != FUNCTION_DECL)
3490 467 : continue;
3491 :
3492 : /* Ignore this variant if its selector is known not to match. */
3493 2281 : if (!omp_context_selector_matches (selector, augmented_context,
3494 : complete_p))
3495 467 : continue;
3496 :
3497 1814 : struct omp_variant candidate;
3498 1814 : candidate.selector = selector;
3499 1814 : candidate.dynamic_selector = false;
3500 1814 : candidate.alternative = fndecl;
3501 1814 : candidate.body = NULL_TREE;
3502 1814 : candidates.safe_push (candidate);
3503 : }
3504 :
3505 : /* Add a default that is the base function. */
3506 2039 : struct omp_variant v;
3507 2039 : v.selector = NULL_TREE;
3508 2039 : v.dynamic_selector = false;
3509 2039 : v.alternative = base;
3510 2039 : v.body = NULL_TREE;
3511 2039 : candidates.safe_push (v);
3512 2039 : return candidates.copy ();
3513 2039 : }
3514 :
3515 : /* Collect metadirective candidates for METADIRECTIVE. CONSTRUCT_CONTEXT
3516 : is the un-augmented context, or NULL_TREE if that information is not
3517 : available yet. */
3518 : vec<struct omp_variant>
3519 494 : omp_metadirective_candidates (tree metadirective, tree construct_context)
3520 : {
3521 494 : auto_vec <struct omp_variant> candidates;
3522 494 : tree variant = OMP_METADIRECTIVE_VARIANTS (metadirective);
3523 494 : bool complete_p;
3524 494 : tree augmented_context
3525 494 : = omp_complete_construct_context (construct_context, &complete_p);
3526 :
3527 494 : gcc_assert (variant);
3528 1578 : for (; variant; variant = TREE_CHAIN (variant))
3529 : {
3530 1084 : tree selector = OMP_METADIRECTIVE_VARIANT_SELECTOR (variant);
3531 :
3532 : /* Ignore this variant if its selector is known not to match. */
3533 1084 : if (!omp_context_selector_matches (selector, augmented_context,
3534 : complete_p))
3535 59 : continue;
3536 :
3537 1025 : struct omp_variant candidate;
3538 1025 : candidate.selector = selector;
3539 1025 : candidate.dynamic_selector = false;
3540 1025 : candidate.alternative = OMP_METADIRECTIVE_VARIANT_DIRECTIVE (variant);
3541 1025 : candidate.body = OMP_METADIRECTIVE_VARIANT_BODY (variant);
3542 1025 : candidates.safe_push (candidate);
3543 : }
3544 494 : return candidates.copy ();
3545 494 : }
3546 :
3547 : /* Return a vector of dynamic replacement candidates for the metadirective
3548 : statement in METADIRECTIVE. Return an empty vector if the metadirective
3549 : cannot be resolved. This function is intended to be called from the
3550 : front ends, prior to gimplification. */
3551 :
3552 : vec<struct omp_variant>
3553 320 : omp_early_resolve_metadirective (tree metadirective)
3554 : {
3555 320 : vec <struct omp_variant> candidates
3556 320 : = omp_metadirective_candidates (metadirective, NULL_TREE);
3557 320 : return omp_get_dynamic_candidates (candidates, NULL_TREE);
3558 : }
3559 :
3560 : /* Return a vector of dynamic replacement candidates for the variant construct
3561 : with SELECTORS and CONSTRUCT_CONTEXT. This version is called during late
3562 : resolution in the ompdevlow pass. */
3563 :
3564 : vec<struct omp_variant>
3565 304 : omp_resolve_variant_construct (tree construct_context, tree selectors)
3566 : {
3567 304 : auto_vec <struct omp_variant> variants;
3568 :
3569 2108 : for (int i = 0; i < TREE_VEC_LENGTH (selectors); i++)
3570 : {
3571 1804 : struct omp_variant variant;
3572 :
3573 1804 : variant.selector = TREE_VEC_ELT (selectors, i);
3574 1804 : variant.dynamic_selector = false;
3575 1804 : variant.alternative = build_int_cst (integer_type_node, i + 1);
3576 1804 : variant.body = NULL_TREE;
3577 :
3578 1804 : variants.safe_push (variant);
3579 : }
3580 :
3581 304 : return omp_get_dynamic_candidates (variants, construct_context);
3582 304 : }
3583 :
3584 : /* Encode an oacc launch argument. This matches the GOMP_LAUNCH_PACK
3585 : macro on gomp-constants.h. We do not check for overflow. */
3586 :
3587 : tree
3588 11070 : oacc_launch_pack (unsigned code, tree device, unsigned op)
3589 : {
3590 11070 : tree res;
3591 :
3592 11070 : res = build_int_cst (unsigned_type_node, GOMP_LAUNCH_PACK (code, 0, op));
3593 11070 : if (device)
3594 : {
3595 0 : device = fold_build2 (LSHIFT_EXPR, unsigned_type_node,
3596 : device, build_int_cst (unsigned_type_node,
3597 : GOMP_LAUNCH_DEVICE_SHIFT));
3598 0 : res = fold_build2 (BIT_IOR_EXPR, unsigned_type_node, res, device);
3599 : }
3600 11070 : return res;
3601 : }
3602 :
3603 : /* Openacc compute grid dimension clauses are converted to an attribute
3604 : attached to the function. This permits the target-side code to (a) massage
3605 : the dimensions, (b) emit that data and (c) optimize. Non-constant
3606 : dimensions are pushed onto ARGS.
3607 :
3608 : The attribute value is a TREE_LIST. A set of dimensions is
3609 : represented as a list of INTEGER_CST. Those that are runtime
3610 : exprs are represented as an INTEGER_CST of zero.
3611 :
3612 : TODO: Normally the attribute will just contain a single such list. If
3613 : however it contains a list of lists, this will represent the use of
3614 : device_type. Each member of the outer list is an assoc list of
3615 : dimensions, keyed by the device type. The first entry will be the
3616 : default. Well, that's the plan. */
3617 :
3618 : /* Replace any existing oacc fn attribute in ATTRIBS with updated
3619 : dimensions. */
3620 :
3621 : tree
3622 20654 : oacc_replace_fn_attrib_attr (tree attribs, tree dims)
3623 : {
3624 20654 : tree ident = get_identifier (OACC_FN_ATTRIB);
3625 :
3626 : /* If we happen to be present as the first attrib, drop it. */
3627 40239 : if (attribs && TREE_PURPOSE (attribs) == ident)
3628 9583 : attribs = TREE_CHAIN (attribs);
3629 20654 : return tree_cons (ident, dims, attribs);
3630 : }
3631 :
3632 : /* Replace any existing oacc fn attribute on FN with updated
3633 : dimensions. */
3634 :
3635 : void
3636 20299 : oacc_replace_fn_attrib (tree fn, tree dims)
3637 : {
3638 20299 : DECL_ATTRIBUTES (fn)
3639 20299 : = oacc_replace_fn_attrib_attr (DECL_ATTRIBUTES (fn), dims);
3640 20299 : }
3641 :
3642 : /* Scan CLAUSES for launch dimensions and attach them to the oacc
3643 : function attribute. Push any that are non-constant onto the ARGS
3644 : list, along with an appropriate GOMP_LAUNCH_DIM tag. */
3645 :
3646 : void
3647 9904 : oacc_set_fn_attrib (tree fn, tree clauses, vec<tree> *args)
3648 : {
3649 : /* Must match GOMP_DIM ordering. */
3650 9904 : static const omp_clause_code ids[]
3651 : = { OMP_CLAUSE_NUM_GANGS, OMP_CLAUSE_NUM_WORKERS,
3652 : OMP_CLAUSE_VECTOR_LENGTH };
3653 9904 : unsigned ix;
3654 9904 : tree dims[GOMP_DIM_MAX];
3655 :
3656 9904 : tree attr = NULL_TREE;
3657 9904 : unsigned non_const = 0;
3658 :
3659 39616 : for (ix = GOMP_DIM_MAX; ix--;)
3660 : {
3661 29712 : tree clause = omp_find_clause (clauses, ids[ix]);
3662 29712 : tree dim = NULL_TREE;
3663 :
3664 29712 : if (clause)
3665 3275 : dim = OMP_CLAUSE_EXPR (clause, ids[ix]);
3666 29712 : dims[ix] = dim;
3667 29712 : if (dim && TREE_CODE (dim) != INTEGER_CST)
3668 : {
3669 157 : dim = integer_zero_node;
3670 157 : non_const |= GOMP_DIM_MASK (ix);
3671 : }
3672 29712 : attr = tree_cons (NULL_TREE, dim, attr);
3673 : }
3674 :
3675 9904 : oacc_replace_fn_attrib (fn, attr);
3676 :
3677 9904 : if (non_const)
3678 : {
3679 : /* Push a dynamic argument set. */
3680 101 : args->safe_push (oacc_launch_pack (GOMP_LAUNCH_DIM,
3681 : NULL_TREE, non_const));
3682 404 : for (unsigned ix = 0; ix != GOMP_DIM_MAX; ix++)
3683 303 : if (non_const & GOMP_DIM_MASK (ix))
3684 157 : args->safe_push (dims[ix]);
3685 : }
3686 9904 : }
3687 :
3688 : /* Verify OpenACC routine clauses.
3689 :
3690 : Returns 0 if FNDECL should be marked with an OpenACC 'routine' directive, 1
3691 : if it has already been marked in compatible way, and -1 if incompatible.
3692 : Upon returning, the chain of clauses will contain exactly one clause
3693 : specifying the level of parallelism. */
3694 :
3695 : int
3696 1219 : oacc_verify_routine_clauses (tree fndecl, tree *clauses, location_t loc,
3697 : const char *routine_str)
3698 : {
3699 1219 : tree c_level = NULL_TREE;
3700 1219 : tree c_nohost = NULL_TREE;
3701 1219 : tree c_p = NULL_TREE;
3702 3627 : for (tree c = *clauses; c; c_p = c, c = OMP_CLAUSE_CHAIN (c))
3703 2408 : switch (OMP_CLAUSE_CODE (c))
3704 : {
3705 2309 : case OMP_CLAUSE_GANG:
3706 2309 : case OMP_CLAUSE_WORKER:
3707 2309 : case OMP_CLAUSE_VECTOR:
3708 2309 : case OMP_CLAUSE_SEQ:
3709 2309 : if (c_level == NULL_TREE)
3710 : c_level = c;
3711 1496 : else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_CODE (c_level))
3712 : {
3713 : /* This has already been diagnosed in the front ends. */
3714 : /* Drop the duplicate clause. */
3715 352 : gcc_checking_assert (c_p != NULL_TREE);
3716 352 : OMP_CLAUSE_CHAIN (c_p) = OMP_CLAUSE_CHAIN (c);
3717 352 : c = c_p;
3718 : }
3719 : else
3720 : {
3721 1144 : error_at (OMP_CLAUSE_LOCATION (c),
3722 : "%qs specifies a conflicting level of parallelism",
3723 1144 : omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
3724 1144 : inform (OMP_CLAUSE_LOCATION (c_level),
3725 : "... to the previous %qs clause here",
3726 1144 : omp_clause_code_name[OMP_CLAUSE_CODE (c_level)]);
3727 : /* Drop the conflicting clause. */
3728 1144 : gcc_checking_assert (c_p != NULL_TREE);
3729 1144 : OMP_CLAUSE_CHAIN (c_p) = OMP_CLAUSE_CHAIN (c);
3730 1144 : c = c_p;
3731 : }
3732 : break;
3733 : case OMP_CLAUSE_NOHOST:
3734 : /* Don't worry about duplicate clauses here. */
3735 : c_nohost = c;
3736 : break;
3737 0 : default:
3738 0 : gcc_unreachable ();
3739 : }
3740 1219 : if (c_level == NULL_TREE)
3741 : {
3742 : /* Default to an implicit 'seq' clause. */
3743 406 : c_level = build_omp_clause (loc, OMP_CLAUSE_SEQ);
3744 406 : OMP_CLAUSE_CHAIN (c_level) = *clauses;
3745 406 : *clauses = c_level;
3746 : }
3747 : /* In *clauses, we now have exactly one clause specifying the level of
3748 : parallelism. */
3749 :
3750 1219 : tree attr
3751 1219 : = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (fndecl));
3752 1219 : if (attr != NULL_TREE)
3753 : {
3754 : /* Diagnose if "#pragma omp declare target" has also been applied. */
3755 429 : if (TREE_VALUE (attr) == NULL_TREE)
3756 : {
3757 : /* See <https://gcc.gnu.org/PR93465>; the semantics of combining
3758 : OpenACC and OpenMP 'target' are not clear. */
3759 32 : error_at (loc,
3760 : "cannot apply %qs to %qD, which has also been"
3761 : " marked with an OpenMP 'declare target' directive",
3762 : routine_str, fndecl);
3763 : /* Incompatible. */
3764 32 : return -1;
3765 : }
3766 :
3767 : /* If a "#pragma acc routine" has already been applied, just verify
3768 : this one for compatibility. */
3769 : /* Collect previous directive's clauses. */
3770 : tree c_level_p = NULL_TREE;
3771 : tree c_nohost_p = NULL_TREE;
3772 850 : for (tree c = TREE_VALUE (attr); c; c = OMP_CLAUSE_CHAIN (c))
3773 453 : switch (OMP_CLAUSE_CODE (c))
3774 : {
3775 397 : case OMP_CLAUSE_GANG:
3776 397 : case OMP_CLAUSE_WORKER:
3777 397 : case OMP_CLAUSE_VECTOR:
3778 397 : case OMP_CLAUSE_SEQ:
3779 397 : gcc_checking_assert (c_level_p == NULL_TREE);
3780 : c_level_p = c;
3781 : break;
3782 56 : case OMP_CLAUSE_NOHOST:
3783 56 : gcc_checking_assert (c_nohost_p == NULL_TREE);
3784 : c_nohost_p = c;
3785 : break;
3786 0 : default:
3787 0 : gcc_unreachable ();
3788 : }
3789 397 : gcc_checking_assert (c_level_p != NULL_TREE);
3790 : /* ..., and compare to current directive's, which we've already collected
3791 : above. */
3792 397 : tree c_diag;
3793 397 : tree c_diag_p;
3794 : /* Matching level of parallelism? */
3795 397 : if (OMP_CLAUSE_CODE (c_level) != OMP_CLAUSE_CODE (c_level_p))
3796 : {
3797 97 : c_diag = c_level;
3798 97 : c_diag_p = c_level_p;
3799 97 : goto incompatible;
3800 : }
3801 : /* Matching 'nohost' clauses? */
3802 300 : if ((c_nohost == NULL_TREE) != (c_nohost_p == NULL_TREE))
3803 : {
3804 56 : c_diag = c_nohost;
3805 56 : c_diag_p = c_nohost_p;
3806 56 : goto incompatible;
3807 : }
3808 : /* Compatible. */
3809 : return 1;
3810 :
3811 153 : incompatible:
3812 153 : if (c_diag != NULL_TREE)
3813 125 : error_at (OMP_CLAUSE_LOCATION (c_diag),
3814 : "incompatible %qs clause when applying"
3815 : " %qs to %qD, which has already been"
3816 : " marked with an OpenACC 'routine' directive",
3817 125 : omp_clause_code_name[OMP_CLAUSE_CODE (c_diag)],
3818 : routine_str, fndecl);
3819 28 : else if (c_diag_p != NULL_TREE)
3820 28 : error_at (loc,
3821 : "missing %qs clause when applying"
3822 : " %qs to %qD, which has already been"
3823 : " marked with an OpenACC 'routine' directive",
3824 28 : omp_clause_code_name[OMP_CLAUSE_CODE (c_diag_p)],
3825 : routine_str, fndecl);
3826 : else
3827 0 : gcc_unreachable ();
3828 153 : if (c_diag_p != NULL_TREE)
3829 125 : inform (OMP_CLAUSE_LOCATION (c_diag_p),
3830 : "... with %qs clause here",
3831 125 : omp_clause_code_name[OMP_CLAUSE_CODE (c_diag_p)]);
3832 : else
3833 : {
3834 : /* In the front ends, we don't preserve location information for the
3835 : OpenACC routine directive itself. However, that of c_level_p
3836 : should be close. */
3837 28 : location_t loc_routine = OMP_CLAUSE_LOCATION (c_level_p);
3838 28 : inform (loc_routine, "... without %qs clause near to here",
3839 28 : omp_clause_code_name[OMP_CLAUSE_CODE (c_diag)]);
3840 : }
3841 : /* Incompatible. */
3842 153 : return -1;
3843 : }
3844 :
3845 : return 0;
3846 : }
3847 :
3848 : /* Process the OpenACC 'routine' directive clauses to generate an attribute
3849 : for the level of parallelism. All dimensions have a size of zero
3850 : (dynamic). TREE_PURPOSE is set to indicate whether that dimension
3851 : can have a loop partitioned on it. non-zero indicates
3852 : yes, zero indicates no. By construction once a non-zero has been
3853 : reached, further inner dimensions must also be non-zero. We set
3854 : TREE_VALUE to zero for the dimensions that may be partitioned and
3855 : 1 for the other ones -- if a loop is (erroneously) spawned at
3856 : an outer level, we don't want to try and partition it. */
3857 :
3858 : tree
3859 1137 : oacc_build_routine_dims (tree clauses)
3860 : {
3861 : /* Must match GOMP_DIM ordering. */
3862 1137 : static const omp_clause_code ids[]
3863 : = {OMP_CLAUSE_GANG, OMP_CLAUSE_WORKER, OMP_CLAUSE_VECTOR, OMP_CLAUSE_SEQ};
3864 1137 : int ix;
3865 1137 : int level = -1;
3866 :
3867 2317 : for (; clauses; clauses = OMP_CLAUSE_CHAIN (clauses))
3868 2301 : for (ix = GOMP_DIM_MAX + 1; ix--;)
3869 2258 : if (OMP_CLAUSE_CODE (clauses) == ids[ix])
3870 : {
3871 : level = ix;
3872 : break;
3873 : }
3874 1137 : gcc_checking_assert (level >= 0);
3875 :
3876 : tree dims = NULL_TREE;
3877 :
3878 4548 : for (ix = GOMP_DIM_MAX; ix--;)
3879 3411 : dims = tree_cons (build_int_cst (boolean_type_node, ix >= level),
3880 3411 : build_int_cst (integer_type_node, ix < level), dims);
3881 :
3882 1137 : return dims;
3883 : }
3884 :
3885 : /* Retrieve the oacc function attrib and return it. Non-oacc
3886 : functions will return NULL. */
3887 :
3888 : tree
3889 351987 : oacc_get_fn_attrib (tree fn)
3890 : {
3891 351987 : return lookup_attribute (OACC_FN_ATTRIB, DECL_ATTRIBUTES (fn));
3892 : }
3893 :
3894 : /* Return true if FN is an OpenMP or OpenACC offloading function. */
3895 :
3896 : bool
3897 561 : offloading_function_p (tree fn)
3898 : {
3899 561 : tree attrs = DECL_ATTRIBUTES (fn);
3900 561 : return (lookup_attribute ("omp declare target", attrs)
3901 561 : || lookup_attribute ("omp target entrypoint", attrs));
3902 : }
3903 :
3904 : /* Extract an oacc execution dimension from FN. FN must be an
3905 : offloaded function or routine that has already had its execution
3906 : dimensions lowered to the target-specific values. */
3907 :
3908 : int
3909 62574 : oacc_get_fn_dim_size (tree fn, int axis)
3910 : {
3911 62574 : tree attrs = oacc_get_fn_attrib (fn);
3912 :
3913 62574 : gcc_assert (axis < GOMP_DIM_MAX);
3914 :
3915 62574 : tree dims = TREE_VALUE (attrs);
3916 127135 : while (axis--)
3917 64561 : dims = TREE_CHAIN (dims);
3918 :
3919 62574 : int size = TREE_INT_CST_LOW (TREE_VALUE (dims));
3920 :
3921 62574 : return size;
3922 : }
3923 :
3924 : /* Extract the dimension axis from an IFN_GOACC_DIM_POS or
3925 : IFN_GOACC_DIM_SIZE call. */
3926 :
3927 : int
3928 93969 : oacc_get_ifn_dim_arg (const gimple *stmt)
3929 : {
3930 93969 : gcc_checking_assert (gimple_call_internal_fn (stmt) == IFN_GOACC_DIM_SIZE
3931 : || gimple_call_internal_fn (stmt) == IFN_GOACC_DIM_POS);
3932 93969 : tree arg = gimple_call_arg (stmt, 0);
3933 93969 : HOST_WIDE_INT axis = TREE_INT_CST_LOW (arg);
3934 :
3935 93969 : gcc_checking_assert (axis >= 0 && axis < GOMP_DIM_MAX);
3936 93969 : return (int) axis;
3937 : }
3938 :
3939 : /* Build COMPONENT_REF and set TREE_THIS_VOLATILE and TREE_READONLY on it
3940 : as appropriate. */
3941 :
3942 : tree
3943 294934 : omp_build_component_ref (tree obj, tree field)
3944 : {
3945 294934 : tree ret = build3 (COMPONENT_REF, TREE_TYPE (field), obj, field, NULL);
3946 294934 : if (TREE_THIS_VOLATILE (field))
3947 62 : TREE_THIS_VOLATILE (ret) |= 1;
3948 294934 : if (TREE_READONLY (field))
3949 0 : TREE_READONLY (ret) |= 1;
3950 294934 : return ret;
3951 : }
3952 :
3953 : /* Return true if NAME is the name of an omp_* runtime API call. */
3954 : bool
3955 8216 : omp_runtime_api_procname (const char *name)
3956 : {
3957 8216 : if (!startswith (name, "omp_"))
3958 : return false;
3959 :
3960 : static const char *omp_runtime_apis[] =
3961 : {
3962 : /* This array has 3 sections. First omp_* calls that don't
3963 : have any suffixes; this implies bind(C) either in the
3964 : specification or implementation choice. */
3965 : "aligned_alloc",
3966 : "aligned_calloc",
3967 : "alloc",
3968 : "calloc",
3969 : "free",
3970 : "get_device_distances",
3971 : "get_interop_int",
3972 : "get_interop_ptr",
3973 : "get_mapped_ptr",
3974 : "get_num_interop_properties",
3975 : "get_num_teams_dim",
3976 : "get_num_threads_dim",
3977 : "get_supported_teams_dim",
3978 : "get_supported_threads_dim",
3979 : "get_team_num_dim",
3980 : "get_thread_limit_dim",
3981 : "get_thread_num_dim",
3982 : "realloc",
3983 : "target_alloc",
3984 : "target_associate_ptr",
3985 : "target_disassociate_ptr",
3986 : "target_free",
3987 : "target_is_accessible",
3988 : "target_is_present",
3989 : "target_memcpy",
3990 : "target_memcpy_async",
3991 : "target_memcpy_rect",
3992 : "target_memcpy_rect_async",
3993 : NULL,
3994 : /* Now omp_* calls that are available as omp_* and omp_*_; however, the
3995 : DECL_NAME is always omp_* without tailing underscore. */
3996 : "capture_affinity",
3997 : "destroy_allocator",
3998 : "destroy_lock",
3999 : "destroy_nest_lock",
4000 : "display_affinity",
4001 : "fulfill_event",
4002 : "get_active_level",
4003 : "get_affinity_format",
4004 : "get_cancellation",
4005 : "get_default_allocator",
4006 : "get_default_device",
4007 : "get_device_from_uid",
4008 : "get_device_num",
4009 : "get_dynamic",
4010 : "get_initial_device",
4011 : "get_interop_name",
4012 : "get_interop_rc_desc",
4013 : "get_interop_str",
4014 : "get_interop_type_desc",
4015 : "get_level",
4016 : "get_max_active_levels",
4017 : "get_max_task_priority",
4018 : "get_max_teams",
4019 : "get_max_threads",
4020 : "get_nested",
4021 : "get_num_devices",
4022 : "get_num_places",
4023 : "get_num_procs",
4024 : "get_num_teams",
4025 : "get_num_threads",
4026 : "get_partition_num_places",
4027 : "get_place_num",
4028 : "get_proc_bind",
4029 : "get_supported_active_league_dims",
4030 : "get_supported_active_levels",
4031 : "get_supported_active_team_dims",
4032 : "get_team_num",
4033 : "get_teams_thread_limit",
4034 : "get_thread_limit",
4035 : "get_thread_num",
4036 : "get_wtick",
4037 : "get_wtime",
4038 : "in_explicit_task",
4039 : "in_final",
4040 : "in_parallel",
4041 : "init_lock",
4042 : "init_nest_lock",
4043 : "is_initial_device",
4044 : "pause_resource",
4045 : "pause_resource_all",
4046 : "set_affinity_format",
4047 : "set_default_allocator",
4048 : "set_lock",
4049 : "set_nest_lock",
4050 : "test_lock",
4051 : "test_nest_lock",
4052 : "unset_lock",
4053 : "unset_nest_lock",
4054 : NULL,
4055 : /* And finally calls available as omp_*, omp_*_ and omp_*_8_; however,
4056 : as DECL_NAME only omp_* and omp_*_8 appear. */
4057 : "control_tool",
4058 : "display_env",
4059 : "get_ancestor_thread_num",
4060 : "get_uid_from_device",
4061 : "get_partition_place_nums",
4062 : "get_place_num_procs",
4063 : "get_place_proc_ids",
4064 : "get_schedule",
4065 : "get_team_size",
4066 : "init_allocator",
4067 : "set_default_device",
4068 : "set_dynamic",
4069 : "set_max_active_levels",
4070 : "set_nested",
4071 : "set_num_teams",
4072 : "set_num_threads",
4073 : "set_schedule",
4074 : "set_teams_thread_limit",
4075 : "target_memset"
4076 : };
4077 :
4078 : int mode = 0;
4079 50618 : for (unsigned i = 0; i < ARRAY_SIZE (omp_runtime_apis); i++)
4080 : {
4081 50617 : if (omp_runtime_apis[i] == NULL)
4082 : {
4083 843 : mode++;
4084 843 : continue;
4085 : }
4086 49774 : size_t len = strlen (omp_runtime_apis[i]);
4087 49774 : if (strncmp (name + 4, omp_runtime_apis[i], len) == 0
4088 995 : && (name[4 + len] == '\0'
4089 3 : || (mode > 1 && strcmp (name + 4 + len, "_8") == 0)))
4090 : return true;
4091 : }
4092 : return false;
4093 : }
4094 :
4095 : /* Return true if FNDECL is an omp_* runtime API call. */
4096 :
4097 : bool
4098 9445 : omp_runtime_api_call (const_tree fndecl)
4099 : {
4100 9445 : tree declname = DECL_NAME (fndecl);
4101 9445 : if (!declname
4102 9445 : || (DECL_CONTEXT (fndecl) != NULL_TREE
4103 6869 : && TREE_CODE (DECL_CONTEXT (fndecl)) != TRANSLATION_UNIT_DECL)
4104 17326 : || !TREE_PUBLIC (fndecl))
4105 : return false;
4106 7738 : return omp_runtime_api_procname (IDENTIFIER_POINTER (declname));
4107 : }
4108 :
4109 : /* See "Additional Definitions for the OpenMP API Specification" document;
4110 : associated IDs are 1, 2, ... */
4111 : static const char* omp_interop_fr_str[] = {"cuda", "cuda_driver", "opencl",
4112 : "sycl", "hip", "level_zero", "hsa"};
4113 :
4114 : /* Returns the foreign-runtime ID if found or 0 otherwise. */
4115 :
4116 : char
4117 523 : omp_get_fr_id_from_name (const char *str)
4118 : {
4119 523 : static_assert (GOMP_INTEROP_IFR_LAST == ARRAY_SIZE (omp_interop_fr_str), "");
4120 :
4121 2005 : for (unsigned i = 0; i < ARRAY_SIZE (omp_interop_fr_str); ++i)
4122 1965 : if (!strcmp (str, omp_interop_fr_str[i]))
4123 483 : return i + 1;
4124 : return GOMP_INTEROP_IFR_UNKNOWN;
4125 : }
4126 :
4127 : /* Returns the string value to a foreign-runtime integer value or NULL if value
4128 : is not known. */
4129 :
4130 : const char *
4131 386 : omp_get_name_from_fr_id (int fr_id)
4132 : {
4133 386 : if (fr_id < 1 || fr_id > (int) ARRAY_SIZE (omp_interop_fr_str))
4134 : return "<unknown>";
4135 272 : return omp_interop_fr_str[fr_id-1];
4136 : }
4137 :
4138 : namespace omp_addr_tokenizer {
4139 :
4140 : /* We scan an expression by recursive descent, and build a vector of
4141 : "omp_addr_token *" pointers representing a "parsed" version of the
4142 : expression. The grammar we use is something like this:
4143 :
4144 : expr0::
4145 : expr [section-access]
4146 :
4147 : expr::
4148 : structured-expr access-method
4149 : | array-base access-method
4150 :
4151 : structured-expr::
4152 : structure-base component-selector
4153 :
4154 : arbitrary-expr::
4155 : (anything else)
4156 :
4157 : structure-base::
4158 : DECL access-method
4159 : | structured-expr access-method
4160 : | arbitrary-expr access-method
4161 :
4162 : array-base::
4163 : DECL
4164 : | arbitrary-expr
4165 :
4166 : access-method::
4167 : DIRECT
4168 : | REF
4169 : | POINTER
4170 : | REF_TO_POINTER
4171 : | POINTER_OFFSET
4172 : | REF_TO_POINTER_OFFSET
4173 : | INDEXED_ARRAY
4174 : | INDEXED_REF_TO_ARRAY
4175 : | index-expr
4176 :
4177 : index-expr::
4178 : INDEX_EXPR access-method
4179 :
4180 : component-selector::
4181 : component-selector COMPONENT_REF
4182 : | component-selector ARRAY_REF
4183 : | COMPONENT_REF
4184 :
4185 : This tokenized form is then used both in parsing, for OpenMP clause
4186 : expansion (for C and C++) and in gimplify.cc for sibling-list handling
4187 : (for C, C++ and Fortran). */
4188 :
4189 33925 : omp_addr_token::omp_addr_token (token_type t, tree e)
4190 33925 : : type(t), expr(e)
4191 : {
4192 33925 : }
4193 :
4194 128163 : omp_addr_token::omp_addr_token (access_method_kinds k, tree e)
4195 128163 : : type(ACCESS_METHOD), expr(e)
4196 : {
4197 128163 : u.access_kind = k;
4198 128163 : }
4199 :
4200 96488 : omp_addr_token::omp_addr_token (token_type t, structure_base_kinds k, tree e)
4201 96488 : : type(t), expr(e)
4202 : {
4203 96488 : u.structure_base_kind = k;
4204 96488 : }
4205 :
4206 : static bool
4207 96492 : omp_parse_component_selector (tree *expr0)
4208 : {
4209 96492 : tree expr = *expr0;
4210 96492 : tree last_component = NULL_TREE;
4211 :
4212 96492 : while (TREE_CODE (expr) == COMPONENT_REF
4213 142655 : || TREE_CODE (expr) == ARRAY_REF)
4214 : {
4215 46163 : if (TREE_CODE (expr) == COMPONENT_REF)
4216 42686 : last_component = expr;
4217 :
4218 46163 : expr = TREE_OPERAND (expr, 0);
4219 :
4220 46163 : if (TREE_CODE (TREE_TYPE (expr)) == REFERENCE_TYPE)
4221 : break;
4222 : }
4223 :
4224 96492 : if (!last_component)
4225 : return false;
4226 :
4227 33925 : *expr0 = last_component;
4228 33925 : return true;
4229 : }
4230 :
4231 : /* This handles references that have had convert_from_reference called on
4232 : them, and also those that haven't. */
4233 :
4234 : static bool
4235 167349 : omp_parse_ref (tree *expr0)
4236 : {
4237 167349 : tree expr = *expr0;
4238 :
4239 167349 : if (TREE_CODE (TREE_TYPE (expr)) == REFERENCE_TYPE)
4240 : return true;
4241 163317 : else if ((TREE_CODE (expr) == INDIRECT_REF
4242 120088 : || (TREE_CODE (expr) == MEM_REF
4243 2 : && integer_zerop (TREE_OPERAND (expr, 1))))
4244 163319 : && TREE_CODE (TREE_TYPE (TREE_OPERAND (expr, 0))) == REFERENCE_TYPE)
4245 : {
4246 9417 : *expr0 = TREE_OPERAND (expr, 0);
4247 9417 : return true;
4248 : }
4249 :
4250 : return false;
4251 : }
4252 :
4253 : static bool
4254 118574 : omp_parse_pointer (tree *expr0, bool *has_offset)
4255 : {
4256 118574 : tree expr = *expr0;
4257 :
4258 118574 : *has_offset = false;
4259 :
4260 118574 : if ((TREE_CODE (expr) == INDIRECT_REF
4261 87534 : || (TREE_CODE (expr) == MEM_REF
4262 1 : && integer_zerop (TREE_OPERAND (expr, 1))))
4263 118575 : && TREE_CODE (TREE_TYPE (TREE_OPERAND (expr, 0))) == POINTER_TYPE)
4264 : {
4265 31041 : expr = TREE_OPERAND (expr, 0);
4266 :
4267 : /* The Fortran FE sometimes emits a no-op cast here. */
4268 31041 : STRIP_NOPS (expr);
4269 :
4270 34405 : while (1)
4271 : {
4272 34405 : if (TREE_CODE (expr) == COMPOUND_EXPR)
4273 : {
4274 101 : expr = TREE_OPERAND (expr, 1);
4275 101 : STRIP_NOPS (expr);
4276 : }
4277 34304 : else if (TREE_CODE (expr) == SAVE_EXPR)
4278 89 : expr = TREE_OPERAND (expr, 0);
4279 34215 : else if (TREE_CODE (expr) == POINTER_PLUS_EXPR)
4280 : {
4281 3174 : *has_offset = true;
4282 3174 : expr = TREE_OPERAND (expr, 0);
4283 : }
4284 : else
4285 : break;
4286 : }
4287 :
4288 31041 : STRIP_NOPS (expr);
4289 :
4290 31041 : *expr0 = expr;
4291 31041 : return true;
4292 : }
4293 :
4294 : return false;
4295 : }
4296 :
4297 : static void
4298 128163 : omp_parse_access_method (tree *expr0, enum access_method_kinds *kind)
4299 : {
4300 128163 : tree expr = *expr0;
4301 128163 : bool has_offset;
4302 :
4303 128163 : if (omp_parse_ref (&expr))
4304 9589 : *kind = ACCESS_REF;
4305 118574 : else if (omp_parse_pointer (&expr, &has_offset))
4306 : {
4307 31041 : if (omp_parse_ref (&expr))
4308 5574 : *kind = has_offset ? ACCESS_REF_TO_POINTER_OFFSET
4309 : : ACCESS_REF_TO_POINTER;
4310 : else
4311 53334 : *kind = has_offset ? ACCESS_POINTER_OFFSET : ACCESS_POINTER;
4312 : }
4313 87533 : else if (TREE_CODE (expr) == ARRAY_REF)
4314 : {
4315 16590 : while (TREE_CODE (expr) == ARRAY_REF)
4316 8445 : expr = TREE_OPERAND (expr, 0);
4317 8145 : if (omp_parse_ref (&expr))
4318 689 : *kind = ACCESS_INDEXED_REF_TO_ARRAY;
4319 : else
4320 7456 : *kind = ACCESS_INDEXED_ARRAY;
4321 : }
4322 : else
4323 79388 : *kind = ACCESS_DIRECT;
4324 :
4325 128163 : STRIP_NOPS (expr);
4326 :
4327 128163 : *expr0 = expr;
4328 128163 : }
4329 :
4330 : static void
4331 128163 : omp_parse_access_methods (vec<omp_addr_token *> &addr_tokens, tree *expr0)
4332 : {
4333 128163 : tree expr = *expr0;
4334 128163 : enum access_method_kinds kind;
4335 128163 : tree am_expr;
4336 :
4337 128163 : omp_parse_access_method (&expr, &kind);
4338 128163 : am_expr = expr;
4339 :
4340 128163 : if (TREE_CODE (expr) == INDIRECT_REF
4341 128163 : || TREE_CODE (expr) == MEM_REF
4342 125390 : || TREE_CODE (expr) == ARRAY_REF)
4343 3172 : omp_parse_access_methods (addr_tokens, &expr);
4344 :
4345 128163 : addr_tokens.safe_push (new omp_addr_token (kind, am_expr));
4346 :
4347 128163 : *expr0 = expr;
4348 128163 : }
4349 :
4350 : static bool omp_parse_structured_expr (vec<omp_addr_token *> &, tree *);
4351 :
4352 : static bool
4353 96488 : omp_parse_structure_base (vec<omp_addr_token *> &addr_tokens,
4354 : tree *expr0, structure_base_kinds *kind,
4355 : vec<omp_addr_token *> &base_access_tokens,
4356 : bool allow_structured = true)
4357 : {
4358 96488 : tree expr = *expr0;
4359 :
4360 96488 : if (allow_structured)
4361 33925 : omp_parse_access_methods (base_access_tokens, &expr);
4362 :
4363 96488 : if (DECL_P (expr))
4364 : {
4365 90491 : *kind = BASE_DECL;
4366 90491 : return true;
4367 : }
4368 :
4369 5997 : if (allow_structured && omp_parse_structured_expr (addr_tokens, &expr))
4370 : {
4371 5422 : *kind = BASE_COMPONENT_EXPR;
4372 5422 : *expr0 = expr;
4373 5422 : return true;
4374 : }
4375 :
4376 575 : *kind = BASE_ARBITRARY_EXPR;
4377 575 : *expr0 = expr;
4378 575 : return true;
4379 : }
4380 :
4381 : static bool
4382 96492 : omp_parse_structured_expr (vec<omp_addr_token *> &addr_tokens, tree *expr0)
4383 : {
4384 96492 : tree expr = *expr0;
4385 96492 : tree base_component = NULL_TREE;
4386 96492 : structure_base_kinds struct_base_kind;
4387 96492 : auto_vec<omp_addr_token *> base_access_tokens;
4388 :
4389 96492 : if (omp_parse_component_selector (&expr))
4390 33925 : base_component = expr;
4391 : else
4392 : return false;
4393 :
4394 33925 : gcc_assert (TREE_CODE (expr) == COMPONENT_REF);
4395 33925 : expr = TREE_OPERAND (expr, 0);
4396 :
4397 33925 : tree structure_base = expr;
4398 :
4399 33925 : if (!omp_parse_structure_base (addr_tokens, &expr, &struct_base_kind,
4400 : base_access_tokens))
4401 : return false;
4402 :
4403 33925 : addr_tokens.safe_push (new omp_addr_token (STRUCTURE_BASE, struct_base_kind,
4404 33925 : structure_base));
4405 33925 : addr_tokens.safe_splice (base_access_tokens);
4406 33925 : addr_tokens.safe_push (new omp_addr_token (COMPONENT_SELECTOR,
4407 33925 : base_component));
4408 :
4409 33925 : *expr0 = expr;
4410 :
4411 33925 : return true;
4412 96492 : }
4413 :
4414 : static bool
4415 62563 : omp_parse_array_expr (vec<omp_addr_token *> &addr_tokens, tree *expr0)
4416 : {
4417 62563 : tree expr = *expr0;
4418 62563 : structure_base_kinds s_kind;
4419 62563 : auto_vec<omp_addr_token *> base_access_tokens;
4420 :
4421 62563 : if (!omp_parse_structure_base (addr_tokens, &expr, &s_kind,
4422 : base_access_tokens, false))
4423 : return false;
4424 :
4425 62563 : addr_tokens.safe_push (new omp_addr_token (ARRAY_BASE, s_kind, expr));
4426 62563 : addr_tokens.safe_splice (base_access_tokens);
4427 :
4428 62563 : *expr0 = expr;
4429 62563 : return true;
4430 62563 : }
4431 :
4432 : /* Return TRUE if the ACCESS_METHOD token at index 'i' has a further
4433 : ACCESS_METHOD chained after it (e.g., if we're processing an expression
4434 : containing multiple pointer indirections). */
4435 :
4436 : bool
4437 48555 : omp_access_chain_p (vec<omp_addr_token *> &addr_tokens, unsigned i)
4438 : {
4439 48555 : gcc_assert (addr_tokens[i]->type == ACCESS_METHOD);
4440 48555 : return (i + 1 < addr_tokens.length ()
4441 48555 : && addr_tokens[i + 1]->type == ACCESS_METHOD);
4442 : }
4443 :
4444 : /* Return the address of the object accessed by the ACCESS_METHOD token
4445 : at 'i': either of the next access method's expr, or of EXPR if we're at
4446 : the end of the list of tokens. */
4447 :
4448 : tree
4449 4591 : omp_accessed_addr (vec<omp_addr_token *> &addr_tokens, unsigned i, tree expr)
4450 : {
4451 4591 : if (i + 1 < addr_tokens.length ())
4452 68 : return build_fold_addr_expr (addr_tokens[i + 1]->expr);
4453 : else
4454 4523 : return build_fold_addr_expr (expr);
4455 : }
4456 :
4457 : } /* namespace omp_addr_tokenizer. */
4458 :
4459 : bool
4460 91066 : omp_parse_expr (vec<omp_addr_token *> &addr_tokens, tree expr)
4461 : {
4462 91066 : using namespace omp_addr_tokenizer;
4463 91066 : auto_vec<omp_addr_token *> expr_access_tokens;
4464 :
4465 91066 : omp_parse_access_methods (expr_access_tokens, &expr);
4466 :
4467 91066 : if (omp_parse_structured_expr (addr_tokens, &expr))
4468 : ;
4469 62563 : else if (omp_parse_array_expr (addr_tokens, &expr))
4470 : ;
4471 : else
4472 : return false;
4473 :
4474 91066 : addr_tokens.safe_splice (expr_access_tokens);
4475 :
4476 91066 : return true;
4477 91066 : }
4478 :
4479 : DEBUG_FUNCTION void
4480 0 : debug_omp_tokenized_addr (vec<omp_addr_token *> &addr_tokens,
4481 : bool with_exprs)
4482 : {
4483 0 : using namespace omp_addr_tokenizer;
4484 0 : const char *sep = with_exprs ? " " : "";
4485 :
4486 0 : for (auto e : addr_tokens)
4487 : {
4488 0 : const char *pfx = "";
4489 :
4490 0 : fputs (sep, stderr);
4491 :
4492 0 : switch (e->type)
4493 : {
4494 0 : case COMPONENT_SELECTOR:
4495 0 : fputs ("component_selector", stderr);
4496 0 : break;
4497 0 : case ACCESS_METHOD:
4498 0 : switch (e->u.access_kind)
4499 : {
4500 0 : case ACCESS_DIRECT:
4501 0 : fputs ("access_direct", stderr);
4502 0 : break;
4503 0 : case ACCESS_REF:
4504 0 : fputs ("access_ref", stderr);
4505 0 : break;
4506 0 : case ACCESS_POINTER:
4507 0 : fputs ("access_pointer", stderr);
4508 0 : break;
4509 0 : case ACCESS_POINTER_OFFSET:
4510 0 : fputs ("access_pointer_offset", stderr);
4511 0 : break;
4512 0 : case ACCESS_REF_TO_POINTER:
4513 0 : fputs ("access_ref_to_pointer", stderr);
4514 0 : break;
4515 0 : case ACCESS_REF_TO_POINTER_OFFSET:
4516 0 : fputs ("access_ref_to_pointer_offset", stderr);
4517 0 : break;
4518 0 : case ACCESS_INDEXED_ARRAY:
4519 0 : fputs ("access_indexed_array", stderr);
4520 0 : break;
4521 0 : case ACCESS_INDEXED_REF_TO_ARRAY:
4522 0 : fputs ("access_indexed_ref_to_array", stderr);
4523 0 : break;
4524 : }
4525 : break;
4526 0 : case ARRAY_BASE:
4527 0 : case STRUCTURE_BASE:
4528 0 : pfx = e->type == ARRAY_BASE ? "array_" : "struct_";
4529 0 : switch (e->u.structure_base_kind)
4530 : {
4531 0 : case BASE_DECL:
4532 0 : fprintf (stderr, "%sbase_decl", pfx);
4533 0 : break;
4534 0 : case BASE_COMPONENT_EXPR:
4535 0 : fputs ("base_component_expr", stderr);
4536 0 : break;
4537 0 : case BASE_ARBITRARY_EXPR:
4538 0 : fprintf (stderr, "%sbase_arbitrary_expr", pfx);
4539 0 : break;
4540 : }
4541 : break;
4542 : }
4543 0 : if (with_exprs)
4544 : {
4545 0 : fputs (" [", stderr);
4546 0 : print_generic_expr (stderr, e->expr);
4547 0 : fputc (']', stderr);
4548 0 : sep = ",\n ";
4549 : }
4550 : else
4551 : sep = " ";
4552 : }
4553 :
4554 0 : fputs ("\n", stderr);
4555 0 : }
4556 :
4557 : /* Return number of iterations of loop I in FOR_STMT. If PSTEP is non-NULL,
4558 : *PSTEP will be the loop step. */
4559 :
4560 : tree
4561 2632 : omp_loop_number_of_iterations (tree for_stmt, int i, tree *pstep)
4562 : {
4563 2632 : tree t = TREE_VEC_ELT (OMP_FOR_INIT (for_stmt), i);
4564 2632 : gcc_assert (TREE_CODE (t) == MODIFY_EXPR);
4565 2632 : tree decl = TREE_OPERAND (t, 0);
4566 2632 : tree n1 = TREE_OPERAND (t, 1);
4567 2632 : tree type = TREE_TYPE (decl);
4568 2632 : tree cond = TREE_VEC_ELT (OMP_FOR_COND (for_stmt), i);
4569 2632 : gcc_assert (COMPARISON_CLASS_P (cond));
4570 2632 : gcc_assert (TREE_OPERAND (cond, 0) == decl);
4571 2632 : tree_code cond_code = TREE_CODE (cond);
4572 2632 : tree n2 = TREE_OPERAND (cond, 1);
4573 2632 : t = TREE_VEC_ELT (OMP_FOR_INCR (for_stmt), i);
4574 2632 : tree step = NULL_TREE;
4575 2632 : switch (TREE_CODE (t))
4576 : {
4577 1141 : case PREINCREMENT_EXPR:
4578 1141 : case POSTINCREMENT_EXPR:
4579 1141 : gcc_assert (!POINTER_TYPE_P (type));
4580 1141 : gcc_assert (TREE_OPERAND (t, 0) == decl);
4581 1141 : step = build_int_cst (type, 1);
4582 1141 : break;
4583 42 : case PREDECREMENT_EXPR:
4584 42 : case POSTDECREMENT_EXPR:
4585 42 : gcc_assert (!POINTER_TYPE_P (type));
4586 42 : gcc_assert (TREE_OPERAND (t, 0) == decl);
4587 42 : step = build_int_cst (type, -1);
4588 42 : break;
4589 1449 : case MODIFY_EXPR:
4590 1449 : gcc_assert (TREE_OPERAND (t, 0) == decl);
4591 1449 : t = TREE_OPERAND (t, 1);
4592 1449 : switch (TREE_CODE (t))
4593 : {
4594 1374 : case PLUS_EXPR:
4595 1374 : if (TREE_OPERAND (t, 1) == decl)
4596 : {
4597 3 : TREE_OPERAND (t, 1) = TREE_OPERAND (t, 0);
4598 3 : TREE_OPERAND (t, 0) = decl;
4599 : }
4600 : /* FALLTHRU */
4601 : case POINTER_PLUS_EXPR:
4602 : case MINUS_EXPR:
4603 1449 : step = omp_get_for_step_from_incr (EXPR_LOCATION (t), t);
4604 1449 : break;
4605 0 : default:
4606 0 : gcc_unreachable ();
4607 : }
4608 1449 : break;
4609 0 : default:
4610 0 : gcc_unreachable ();
4611 : }
4612 2632 : omp_adjust_for_condition (EXPR_LOCATION (for_stmt), &cond_code, &n2,
4613 : decl, step);
4614 2632 : if (pstep)
4615 2094 : *pstep = step;
4616 2632 : if (INTEGRAL_TYPE_P (type)
4617 2632 : && TYPE_PRECISION (type) < TYPE_PRECISION (long_long_integer_type_node))
4618 : {
4619 2523 : n1 = fold_convert (long_long_integer_type_node, n1);
4620 2523 : n2 = fold_convert (long_long_integer_type_node, n2);
4621 2523 : step = fold_convert (long_long_integer_type_node, step);
4622 : }
4623 2632 : if (cond_code == LT_EXPR
4624 156 : || POINTER_TYPE_P (type)
4625 2768 : || !TYPE_UNSIGNED (TREE_TYPE (n1)))
4626 : {
4627 2632 : if (POINTER_TYPE_P (type))
4628 60 : t = fold_build2 (POINTER_DIFF_EXPR, ssizetype, n2, n1);
4629 : else
4630 2572 : t = fold_build2 (MINUS_EXPR, TREE_TYPE (n1), n2, n1);
4631 2632 : t = fold_build2 (CEIL_DIV_EXPR, TREE_TYPE (t), t, step);
4632 : }
4633 : else
4634 : {
4635 0 : t = fold_build2 (MINUS_EXPR, type, n1, n2);
4636 0 : t = fold_build2 (CEIL_DIV_EXPR, type, t,
4637 : fold_build1 (NEGATE_EXPR, type, step));
4638 : }
4639 2632 : return t;
4640 : }
4641 :
4642 : /* Tile transformation:
4643 : Original loop:
4644 :
4645 : #pragma omp tile sizes(16, 32)
4646 : for (i = 0; i < k; ++i)
4647 : for (j = 0; j < 128; j += 2)
4648 : {
4649 : baz (i, j);
4650 : }
4651 :
4652 : Transformed loop:
4653 : #pragma omp tile sizes(16, 32)
4654 : for (i.0 = 0; i.0 < k; i.0 += 16)
4655 : for (j.0 = 0; j.0 < 128; j.0 += 64)
4656 : {
4657 : i = i.0;
4658 : i.1 = MIN_EXPR <i.0 + 16, k>;
4659 : goto <D.2783>;
4660 : <D.2782>:;
4661 : j = j.0;
4662 : j.1 = j.0 + 32;
4663 : goto <D.2786>;
4664 : <D.2785>:;
4665 : {
4666 : baz (i, j);
4667 : }
4668 : j += 2;
4669 : <D.2786>:;
4670 : if (j < j.1) goto <D.2785>; else goto <D.2787>;
4671 : <D.2787>:;
4672 : ++i;
4673 : <D.2783>:;
4674 : if (i < i.1) goto <D.2782>; else goto <D.2784>;
4675 : <D.2784>:;
4676 : }
4677 :
4678 : where the grid loops have canonical form, but the inner
4679 : loops don't and so are immediately lowered. */
4680 :
4681 : static void
4682 1953 : omp_apply_tile (tree for_stmt, tree sizes, int size)
4683 : {
4684 1953 : tree pre_body = NULL_TREE, post_body = NULL_TREE;
4685 1953 : tree orig_sizes = sizes;
4686 1953 : if (OMP_FOR_NON_RECTANGULAR (for_stmt))
4687 : {
4688 51 : error_at (EXPR_LOCATION (for_stmt), "non-rectangular %<tile%>");
4689 51 : return;
4690 : }
4691 4424 : for (int i = 0; i < TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt)); i++)
4692 : {
4693 2522 : if (orig_sizes)
4694 : {
4695 1734 : size = tree_to_uhwi (TREE_VALUE (sizes));
4696 1734 : sizes = TREE_CHAIN (sizes);
4697 : }
4698 2522 : if (size == 1)
4699 428 : continue;
4700 2094 : if (OMP_FOR_ORIG_DECLS (for_stmt) == NULL_TREE)
4701 : {
4702 517 : OMP_FOR_ORIG_DECLS (for_stmt)
4703 517 : = make_tree_vec (TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt)));
4704 1178 : for (int j = 0; j < TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt)); j++)
4705 : {
4706 661 : gcc_assert (TREE_CODE (TREE_VEC_ELT (OMP_FOR_INIT (for_stmt), j))
4707 : == MODIFY_EXPR);
4708 661 : TREE_VEC_ELT (OMP_FOR_ORIG_DECLS (for_stmt), j)
4709 1322 : = TREE_OPERAND (TREE_VEC_ELT (OMP_FOR_INIT (for_stmt), j), 0);
4710 : }
4711 : }
4712 2094 : tree step;
4713 2094 : tree iters = omp_loop_number_of_iterations (for_stmt, i, &step);
4714 2094 : tree t = TREE_VEC_ELT (OMP_FOR_INIT (for_stmt), i);
4715 2094 : tree decl = TREE_OPERAND (t, 0);
4716 2094 : tree type = TREE_TYPE (decl);
4717 2094 : tree griddecl = create_tmp_var_raw (type);
4718 2094 : DECL_CONTEXT (griddecl) = current_function_decl;
4719 2094 : t = build1 (DECL_EXPR, void_type_node, griddecl);
4720 2094 : append_to_statement_list (t, &OMP_FOR_PRE_BODY (for_stmt));
4721 2094 : TREE_OPERAND (TREE_VEC_ELT (OMP_FOR_INIT (for_stmt), i), 0) = griddecl;
4722 2094 : TREE_PRIVATE (TREE_VEC_ELT (OMP_FOR_INIT (for_stmt), i)) = 1;
4723 2094 : tree cond = TREE_VEC_ELT (OMP_FOR_COND (for_stmt), i);
4724 2094 : TREE_OPERAND (cond, 0) = griddecl;
4725 2094 : tree ub = save_expr (TREE_OPERAND (cond, 1));
4726 2094 : TREE_OPERAND (cond, 1) = ub;
4727 2094 : t = TREE_VEC_ELT (OMP_FOR_INCR (for_stmt), i);
4728 2094 : if (TREE_CODE (cond) == NE_EXPR)
4729 : {
4730 163 : tree_code cond_code = TREE_CODE (cond);
4731 163 : omp_adjust_for_condition (EXPR_LOCATION (for_stmt), &cond_code,
4732 : &ub, griddecl, step);
4733 163 : TREE_SET_CODE (cond, cond_code);
4734 : }
4735 2094 : step = save_expr (step);
4736 2094 : tree gridstep = fold_build2 (MULT_EXPR, TREE_TYPE (step),
4737 : step, build_int_cst (TREE_TYPE (step),
4738 : size));
4739 2094 : if (POINTER_TYPE_P (type))
4740 52 : t = build2 (POINTER_PLUS_EXPR, type, griddecl,
4741 : fold_convert (sizetype, gridstep));
4742 : else
4743 2042 : t = build2 (PLUS_EXPR, type, griddecl, gridstep);
4744 2094 : t = build2 (MODIFY_EXPR, type, griddecl, t);
4745 2094 : TREE_VEC_ELT (OMP_FOR_INCR (for_stmt), i) = t;
4746 2094 : t = build2 (MODIFY_EXPR, type, decl, griddecl);
4747 2094 : append_to_statement_list (t, &pre_body);
4748 2094 : if (POINTER_TYPE_P (type))
4749 52 : t = build2 (POINTER_PLUS_EXPR, type, griddecl,
4750 : fold_convert (sizetype, gridstep));
4751 : else
4752 2042 : t = build2 (PLUS_EXPR, type, griddecl, gridstep);
4753 2094 : bool minmax_needed = true;
4754 2094 : if (TREE_CODE (iters) == INTEGER_CST)
4755 : {
4756 1130 : wide_int witers = wi::to_wide (iters);
4757 1130 : wide_int wsize = wide_int::from (size, witers.get_precision (),
4758 1130 : TYPE_SIGN (TREE_TYPE (iters)));
4759 1130 : if (wi::multiple_of_p (witers, wsize, TYPE_SIGN (TREE_TYPE (iters))))
4760 412 : minmax_needed = false;
4761 1130 : }
4762 1130 : if (minmax_needed)
4763 1682 : switch (TREE_CODE (cond))
4764 : {
4765 251 : case LE_EXPR:
4766 251 : if (POINTER_TYPE_P (type))
4767 4 : t = build2 (MIN_EXPR, type, t,
4768 : build2 (POINTER_PLUS_EXPR, type, ub, size_int (1)));
4769 : else
4770 247 : t = build2 (MIN_EXPR, type, t,
4771 : build2 (PLUS_EXPR, type, ub, build_one_cst (type)));
4772 : break;
4773 1375 : case LT_EXPR:
4774 1375 : t = build2 (MIN_EXPR, type, t, ub);
4775 1375 : break;
4776 28 : case GE_EXPR:
4777 28 : if (POINTER_TYPE_P (type))
4778 8 : t = build2 (MAX_EXPR, type, t,
4779 : build2 (POINTER_PLUS_EXPR, type, ub, size_int (-1)));
4780 : else
4781 20 : t = build2 (MAX_EXPR, type, t,
4782 : build2 (PLUS_EXPR, type, ub,
4783 : build_minus_one_cst (type)));
4784 : break;
4785 28 : case GT_EXPR:
4786 28 : t = build2 (MAX_EXPR, type, t, ub);
4787 28 : break;
4788 0 : default:
4789 0 : gcc_unreachable ();
4790 : }
4791 2094 : tree end = create_tmp_var_raw (type);
4792 2094 : DECL_CONTEXT (end) = current_function_decl;
4793 2094 : end = build4 (TARGET_EXPR, type, end, t, NULL_TREE, NULL_TREE);
4794 2094 : TREE_SIDE_EFFECTS (end) = 1;
4795 2094 : append_to_statement_list (end, &pre_body);
4796 2094 : tree lab1 = create_artificial_label (UNKNOWN_LOCATION);
4797 2094 : tree lab2 = create_artificial_label (UNKNOWN_LOCATION);
4798 2094 : t = build1 (GOTO_EXPR, void_type_node, lab2);
4799 2094 : append_to_statement_list (t, &pre_body);
4800 2094 : t = build1 (LABEL_EXPR, void_type_node, lab1);
4801 2094 : append_to_statement_list (t, &pre_body);
4802 2094 : tree this_post_body = NULL_TREE;
4803 2094 : if (POINTER_TYPE_P (type))
4804 52 : t = build2 (POINTER_PLUS_EXPR, type, decl,
4805 : fold_convert (sizetype, step));
4806 : else
4807 2042 : t = build2 (PLUS_EXPR, type, decl, step);
4808 2094 : t = build2 (MODIFY_EXPR, type, decl, t);
4809 2094 : append_to_statement_list (t, &this_post_body);
4810 2094 : t = build1 (LABEL_EXPR, void_type_node, lab2);
4811 2094 : append_to_statement_list (t, &this_post_body);
4812 2160 : t = build2 ((TREE_CODE (cond) == LT_EXPR || TREE_CODE (cond) == LE_EXPR)
4813 : ? LT_EXPR : GT_EXPR, boolean_type_node, decl, end);
4814 2094 : if (orig_sizes == NULL_TREE)
4815 : {
4816 740 : gcc_assert (i == 0);
4817 740 : t = build3 (ANNOTATE_EXPR, TREE_TYPE (t), t,
4818 : build_int_cst (integer_type_node,
4819 : annot_expr_unroll_kind),
4820 740 : build_int_cst (integer_type_node, size));
4821 : }
4822 2094 : t = build3 (COND_EXPR, void_type_node, t,
4823 : build1 (GOTO_EXPR, void_type_node, lab1), NULL_TREE);
4824 2094 : append_to_statement_list (t, &this_post_body);
4825 2094 : append_to_statement_list (post_body, &this_post_body);
4826 2094 : post_body = this_post_body;
4827 : }
4828 1902 : if (pre_body || post_body)
4829 : {
4830 1609 : append_to_statement_list (OMP_FOR_BODY (for_stmt), &pre_body);
4831 1609 : append_to_statement_list (post_body, &pre_body);
4832 1609 : OMP_FOR_BODY (for_stmt) = pre_body;
4833 : }
4834 : }
4835 :
4836 : /* Callback for walk_tree to find nested loop transforming construct. */
4837 :
4838 : static tree
4839 9281 : find_nested_loop_xform (tree *tp, int *walk_subtrees, void *data)
4840 : {
4841 9281 : tree **pdata = (tree **) data;
4842 9281 : *walk_subtrees = 0;
4843 9281 : switch (TREE_CODE (*tp))
4844 : {
4845 1982 : case OMP_TILE:
4846 1982 : case OMP_UNROLL:
4847 1982 : pdata[1] = tp;
4848 1982 : return *tp;
4849 2615 : case BIND_EXPR:
4850 2615 : if (BIND_EXPR_VARS (*tp)
4851 2615 : || (BIND_EXPR_BLOCK (*tp)
4852 1304 : && BLOCK_VARS (BIND_EXPR_BLOCK (*tp))))
4853 1311 : pdata[0] = tp;
4854 2615 : *walk_subtrees = 1;
4855 2615 : break;
4856 420 : case STATEMENT_LIST:
4857 420 : if (!tsi_one_before_end_p (tsi_start (*tp)))
4858 420 : pdata[0] = tp;
4859 420 : *walk_subtrees = 1;
4860 420 : break;
4861 32 : case TRY_FINALLY_EXPR:
4862 32 : case CLEANUP_POINT_EXPR:
4863 32 : pdata[0] = tp;
4864 32 : *walk_subtrees = 1;
4865 32 : break;
4866 : default:
4867 : break;
4868 : }
4869 : return NULL;
4870 : }
4871 :
4872 : /* Main entry point for performing OpenMP loop transformations. */
4873 :
4874 : void
4875 60583 : omp_maybe_apply_loop_xforms (tree *expr_p, tree for_clauses)
4876 : {
4877 60583 : tree for_stmt = *expr_p;
4878 :
4879 60583 : switch (TREE_CODE (for_stmt))
4880 : {
4881 3908 : case OMP_TILE:
4882 3908 : case OMP_UNROLL:
4883 3908 : if (OMP_LOOPXFORM_LOWERED (for_stmt))
4884 : return;
4885 : break;
4886 : default:
4887 : break;
4888 : }
4889 :
4890 : tree *inner_expr_p = expr_p;
4891 : tree inner_for_stmt = for_stmt;
4892 143053 : for (int i = 0; i < TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt)); i++)
4893 : {
4894 : /* If some loop nest needs one or more loops in canonical form
4895 : from nested loop transforming constructs, first perform the
4896 : loop transformation on the nested construct and then move over
4897 : the corresponding loops in canonical form from the inner construct
4898 : to the outer one. */
4899 84187 : if (TREE_VEC_ELT (OMP_FOR_INIT (for_stmt), i) == NULL_TREE)
4900 : {
4901 1511 : if (inner_for_stmt == for_stmt
4902 2998 : && omp_find_clause (for_clauses ? for_clauses
4903 1487 : : OMP_FOR_CLAUSES (for_stmt),
4904 : OMP_CLAUSE_ORDERED))
4905 : {
4906 73 : error_at (EXPR_LOCATION (for_stmt),
4907 : "%<ordered%> clause used with generated loops");
4908 73 : *expr_p = void_node;
4909 73 : return;
4910 : }
4911 1438 : tree *data[2] = { NULL, NULL };
4912 1438 : walk_tree (&OMP_FOR_BODY (inner_for_stmt),
4913 : find_nested_loop_xform, &data, NULL);
4914 1438 : gcc_assert (data[1]);
4915 1438 : if (data[0])
4916 : {
4917 : /* If there is a BIND_EXPR declaring some vars, or statement
4918 : list with more than one stmt etc., move the intervening
4919 : code around the outermost loop. */
4920 1003 : tree t = *inner_expr_p;
4921 1003 : *inner_expr_p = OMP_FOR_BODY (inner_for_stmt);
4922 1003 : OMP_FOR_BODY (inner_for_stmt) = *data[1];
4923 1003 : *data[1] = t;
4924 1003 : inner_expr_p = data[1];
4925 1003 : data[1] = &OMP_FOR_BODY (inner_for_stmt);
4926 : }
4927 1438 : inner_for_stmt = *data[1];
4928 :
4929 1438 : omp_maybe_apply_loop_xforms (data[1], NULL_TREE);
4930 1438 : if (*data[1] != inner_for_stmt)
4931 : {
4932 544 : tree *data2[2] = { NULL, NULL };
4933 544 : walk_tree (data[1], find_nested_loop_xform, &data2, NULL);
4934 544 : gcc_assert (data2[1]
4935 : && *data2[1] == inner_for_stmt
4936 : && data2[0]);
4937 544 : tree t = *inner_expr_p;
4938 544 : *inner_expr_p = *data[1];
4939 544 : *data[1] = *data2[1];
4940 544 : *data2[1] = t;
4941 544 : inner_expr_p = data2[1];
4942 : }
4943 1438 : tree clauses = OMP_FOR_CLAUSES (inner_for_stmt);
4944 1438 : gcc_checking_assert (TREE_CODE (inner_for_stmt) != OMP_UNROLL
4945 : || omp_find_clause (clauses,
4946 : OMP_CLAUSE_PARTIAL));
4947 1438 : append_to_statement_list (OMP_FOR_PRE_BODY (inner_for_stmt),
4948 : &OMP_FOR_PRE_BODY (for_stmt));
4949 1438 : OMP_FOR_PRE_BODY (inner_for_stmt) = NULL_TREE;
4950 1438 : if (OMP_FOR_ORIG_DECLS (for_stmt) == NULL_TREE
4951 1438 : && OMP_FOR_ORIG_DECLS (inner_for_stmt) != NULL_TREE)
4952 : {
4953 619 : OMP_FOR_ORIG_DECLS (for_stmt)
4954 619 : = make_tree_vec (TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt)));
4955 1456 : for (int j = 0; j < TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt));
4956 : j++)
4957 : {
4958 837 : if (TREE_VEC_ELT (OMP_FOR_INIT (for_stmt), j) == NULL_TREE)
4959 712 : continue;
4960 125 : gcc_assert (TREE_CODE (TREE_VEC_ELT (OMP_FOR_INIT (for_stmt),
4961 : j)) == MODIFY_EXPR);
4962 125 : TREE_VEC_ELT (OMP_FOR_ORIG_DECLS (for_stmt), j)
4963 250 : = TREE_OPERAND (TREE_VEC_ELT (OMP_FOR_INIT (for_stmt), j),
4964 : 0);
4965 : }
4966 : }
4967 3094 : for (int j = 0; j < TREE_VEC_LENGTH (OMP_FOR_INIT (inner_for_stmt));
4968 : ++j)
4969 : {
4970 1919 : if (i + j == TREE_VEC_LENGTH (OMP_FOR_INIT (for_stmt)))
4971 : break;
4972 1656 : if (OMP_FOR_ORIG_DECLS (for_stmt))
4973 : {
4974 1571 : if (OMP_FOR_ORIG_DECLS (inner_for_stmt))
4975 : {
4976 1571 : TREE_VEC_ELT (OMP_FOR_ORIG_DECLS (for_stmt), i + j)
4977 1571 : = TREE_VEC_ELT (OMP_FOR_ORIG_DECLS (inner_for_stmt),
4978 : j);
4979 1571 : TREE_VEC_ELT (OMP_FOR_ORIG_DECLS (inner_for_stmt), j)
4980 1571 : = NULL_TREE;
4981 : }
4982 : else
4983 : {
4984 0 : tree t = TREE_VEC_ELT (OMP_FOR_INIT (inner_for_stmt), j);
4985 0 : gcc_assert (TREE_CODE (t) == MODIFY_EXPR);
4986 0 : TREE_VEC_ELT (OMP_FOR_ORIG_DECLS (for_stmt), i + j)
4987 0 : = TREE_OPERAND (t, 0);
4988 : }
4989 : }
4990 1656 : TREE_VEC_ELT (OMP_FOR_INIT (for_stmt), i + j)
4991 1656 : = TREE_VEC_ELT (OMP_FOR_INIT (inner_for_stmt), j);
4992 1656 : TREE_VEC_ELT (OMP_FOR_COND (for_stmt), i + j)
4993 1656 : = TREE_VEC_ELT (OMP_FOR_COND (inner_for_stmt), j);
4994 1656 : TREE_VEC_ELT (OMP_FOR_INCR (for_stmt), i + j)
4995 1656 : = TREE_VEC_ELT (OMP_FOR_INCR (inner_for_stmt), j);
4996 1656 : TREE_VEC_ELT (OMP_FOR_INIT (inner_for_stmt), j) = NULL_TREE;
4997 1656 : TREE_VEC_ELT (OMP_FOR_COND (inner_for_stmt), j) = NULL_TREE;
4998 1656 : TREE_VEC_ELT (OMP_FOR_INCR (inner_for_stmt), j) = NULL_TREE;
4999 : }
5000 : }
5001 : }
5002 :
5003 58866 : switch (TREE_CODE (for_stmt))
5004 : {
5005 1165 : case OMP_TILE:
5006 1165 : tree sizes;
5007 1165 : sizes = omp_find_clause (OMP_FOR_CLAUSES (for_stmt), OMP_CLAUSE_SIZES);
5008 1165 : omp_apply_tile (for_stmt, OMP_CLAUSE_SIZES_LIST (sizes), 0);
5009 1165 : OMP_LOOPXFORM_LOWERED (for_stmt) = 1;
5010 1165 : break;
5011 1099 : case OMP_UNROLL:
5012 1099 : tree partial;
5013 1099 : partial = omp_find_clause (OMP_FOR_CLAUSES (for_stmt),
5014 : OMP_CLAUSE_PARTIAL);
5015 1099 : if (partial)
5016 788 : omp_apply_tile (for_stmt, NULL_TREE,
5017 788 : OMP_CLAUSE_PARTIAL_EXPR (partial)
5018 604 : ? tree_to_shwi (OMP_CLAUSE_PARTIAL_EXPR (partial))
5019 : : 8);
5020 311 : else if (omp_find_clause (OMP_FOR_CLAUSES (for_stmt), OMP_CLAUSE_FULL))
5021 : {
5022 269 : tree iters = omp_loop_number_of_iterations (for_stmt, 0, NULL);
5023 269 : if (TREE_CODE (iters) != INTEGER_CST)
5024 12 : error_at (EXPR_LOCATION (for_stmt),
5025 : "non-constant iteration count of %<unroll full%> loop");
5026 : }
5027 1099 : OMP_LOOPXFORM_LOWERED (for_stmt) = 1;
5028 1099 : break;
5029 : default:
5030 : break;
5031 : }
5032 : }
5033 :
5034 : /* The next group of functions support merging of context selectors for
5035 : nested "begin declare variant" directives. The spec says:
5036 :
5037 : ...the effective context selectors of the outer directive are
5038 : appended to the context selector of the inner directive to form the
5039 : effective context selector of the inner directive. If a
5040 : trait-set-selector is present on both directives, the trait-selector
5041 : list of the outer directive is appended to the trait-selector list
5042 : of the inner directive after equivalent trait-selectors have been
5043 : removed from the outer list.
5044 :
5045 : In other words, there is no requirement to combine non-equivalent
5046 : trait-selectors according to their peculiar semantics, such as allowing
5047 : "any" as a wildcard, ANDing trait-property-expressions of "condition" trait
5048 : property expressions together, etc. Also there is no special provision for
5049 : treating the "construct" selector as an ordered list.
5050 :
5051 : Note that the spec does not explicitly say what "equivalent" means;
5052 : whether the properties and score of the trait-selectors must be identical,
5053 : or only the name of the trait-selector. This code assumes the former
5054 : except for the construct trait set where the order of selectors
5055 : is significant (so that it is *not* equivalent to have the same
5056 : trait-selector appearing in a different order in the list). */
5057 :
5058 : /* Copy traits from FROM_TS and push them onto TAIL. */
5059 :
5060 : static tree
5061 36 : omp_copy_trait_set (tree from_ts, tree tail)
5062 : {
5063 92 : for (tree ts = from_ts; ts; ts = TREE_CHAIN (ts))
5064 112 : tail = make_trait_selector (OMP_TS_CODE (ts), OMP_TS_SCORE (ts),
5065 56 : OMP_TS_PROPERTIES (ts), tail);
5066 36 : return nreverse (tail);
5067 : }
5068 :
5069 : /* Return true if trait selectors TS1 and TS2 for set TSS are "equivalent". */
5070 :
5071 : static bool
5072 8 : omp_trait_selectors_equivalent (tree ts1, tree ts2, enum omp_tss_code tss)
5073 : {
5074 8 : if (OMP_TS_CODE (ts1) != OMP_TS_CODE (ts2))
5075 : return false;
5076 :
5077 8 : tree score1 = OMP_TS_SCORE (ts1);
5078 8 : tree score2 = OMP_TS_SCORE (ts2);
5079 0 : if ((score1 && score2 && !simple_cst_equal (score1, score2))
5080 8 : || (score1 && !score2)
5081 16 : || (!score1 && score2))
5082 0 : return false;
5083 :
5084 16 : return (omp_context_selector_props_compare (tss, OMP_TS_CODE (ts1),
5085 8 : OMP_TS_PROPERTIES (ts1),
5086 8 : OMP_TS_PROPERTIES (ts2))
5087 8 : == 0);
5088 : }
5089 :
5090 : /* Merge lists of the trait selectors OUTER_TS and INNER_TS for selector set
5091 : TSS: "the trait-selector list of the outer directive is appended to the
5092 : trait-selector list of the inner directive after equivalent trait-selectors
5093 : have been removed from the outer list". */
5094 :
5095 : static tree
5096 20 : omp_combine_trait_sets (tree outer_ts, tree inner_ts, enum omp_tss_code tss)
5097 : {
5098 20 : unsigned HOST_WIDE_INT inner_traits = 0;
5099 20 : tree to_list = NULL_TREE;
5100 :
5101 40 : for (tree inner = inner_ts; inner; inner = TREE_CHAIN (inner))
5102 : {
5103 20 : omp_ts_code ts_code = OMP_TS_CODE (inner);
5104 20 : inner_traits |= 1 << ts_code;
5105 20 : to_list
5106 20 : = make_trait_selector (ts_code, OMP_TS_SCORE (inner),
5107 20 : unshare_expr (OMP_TS_PROPERTIES (inner)),
5108 : to_list);
5109 : }
5110 :
5111 48 : for (tree outer = outer_ts; outer; outer = TREE_CHAIN (outer))
5112 : {
5113 28 : omp_ts_code ts_code = OMP_TS_CODE (outer);
5114 28 : bool remove = false;
5115 28 : if (inner_traits & (1 << ts_code))
5116 8 : for (tree inner = inner_ts; inner; inner = TREE_CHAIN (inner))
5117 8 : if (OMP_TS_CODE (inner) == ts_code)
5118 : {
5119 8 : if (omp_trait_selectors_equivalent (inner, outer, tss))
5120 : remove = true;
5121 : break;
5122 : }
5123 : if (!remove)
5124 20 : to_list
5125 20 : = make_trait_selector (ts_code, OMP_TS_SCORE (outer),
5126 20 : unshare_expr (OMP_TS_PROPERTIES (outer)),
5127 : to_list);
5128 : }
5129 :
5130 20 : return nreverse (to_list);
5131 : }
5132 :
5133 : /* Merge context selector INNER_CTX with OUTER_CTX. LOC and DIRECTIVE are
5134 : used for error checking. Returns the merged context, or error_mark_node
5135 : if the contexts cannot be merged. */
5136 :
5137 : tree
5138 36 : omp_merge_context_selectors (location_t loc, tree outer_ctx, tree inner_ctx,
5139 : enum omp_ctx_directive directive)
5140 : {
5141 36 : tree merged_ctx = NULL_TREE;
5142 :
5143 36 : if (inner_ctx == error_mark_node || outer_ctx == error_mark_node)
5144 : return error_mark_node;
5145 :
5146 216 : for (unsigned i = OMP_TRAIT_SET_CONSTRUCT; i != OMP_TRAIT_SET_LAST; i++)
5147 : {
5148 180 : omp_tss_code tss_code = static_cast<omp_tss_code>(i);
5149 180 : tree outer_ts = omp_get_context_selector_list (outer_ctx, tss_code);
5150 180 : tree inner_ts = omp_get_context_selector_list (inner_ctx, tss_code);
5151 180 : tree merged_ts = NULL_TREE;
5152 :
5153 180 : if (inner_ts && outer_ts)
5154 20 : merged_ts = omp_combine_trait_sets (outer_ts, inner_ts, tss_code);
5155 160 : else if (inner_ts)
5156 16 : merged_ts = omp_copy_trait_set (inner_ts, NULL_TREE);
5157 144 : else if (outer_ts)
5158 20 : merged_ts = omp_copy_trait_set (outer_ts, NULL_TREE);
5159 :
5160 56 : if (merged_ts)
5161 56 : merged_ctx = make_trait_set_selector (tss_code, merged_ts,
5162 : merged_ctx);
5163 : }
5164 :
5165 36 : merged_ctx = nreverse (merged_ctx);
5166 36 : return omp_check_context_selector (loc, merged_ctx, directive);
5167 : }
5168 :
5169 : /* Remove duplicate and merge clauses mapping the same variable. This function
5170 : is called twice: FIRST in the C and C++ front-ends before any clause
5171 : expansion happens, then in the gimplifier before gathering groups. This is
5172 : because it is easier to process most clauses earlier but some duplicates
5173 : still get introduced during the early clause expansion in the front-ends. */
5174 :
5175 : tree
5176 91746 : omp_remove_duplicate_maps (tree clauses, bool first)
5177 : {
5178 91746 : if (clauses == NULL_TREE)
5179 : return NULL_TREE;
5180 :
5181 70170 : tree outlist = NULL_TREE;
5182 70170 : tree *outlist_p = &outlist;
5183 70170 : bool remove = false;
5184 70170 : tree c1;
5185 235342 : for (c1 = clauses; OMP_CLAUSE_CHAIN (c1) != NULL_TREE;
5186 82586 : c1 = OMP_CLAUSE_CHAIN (c1))
5187 : {
5188 82586 : if (OMP_CLAUSE_CODE (c1) != OMP_CLAUSE_MAP)
5189 : {
5190 58243 : *outlist_p = c1;
5191 58243 : outlist_p = &OMP_CLAUSE_CHAIN (*outlist_p);
5192 58243 : continue;
5193 : }
5194 :
5195 124116 : for (tree c2 = OMP_CLAUSE_CHAIN (c1); c2 != NULL_TREE;
5196 99773 : c2 = OMP_CLAUSE_CHAIN (c2))
5197 : {
5198 99985 : if (OMP_CLAUSE_CODE (c2) != OMP_CLAUSE_MAP)
5199 18198 : continue;
5200 :
5201 81787 : bool maybe_dup_found
5202 81787 : = (OMP_CLAUSE_CODE (c1) == OMP_CLAUSE_CODE (c2)
5203 81787 : && ((/* In the current state, a map clause decl is not supposed
5204 : to be NULL; but let's be defensive. */
5205 81787 : OMP_CLAUSE_DECL (c1) == NULL_TREE
5206 0 : && OMP_CLAUSE_DECL (c2) == NULL_TREE)
5207 81787 : || operand_equal_p (OMP_CLAUSE_DECL (c1),
5208 81787 : OMP_CLAUSE_DECL (c2)))
5209 970 : && ((/* The clause size is generally not known right after
5210 : parsing. */
5211 970 : OMP_CLAUSE_SIZE (c1) == NULL_TREE
5212 308 : && OMP_CLAUSE_SIZE (c2) == NULL_TREE)
5213 673 : || (OMP_CLAUSE_SIZE (c1) != NULL_TREE
5214 662 : && OMP_CLAUSE_SIZE (c2) != NULL_TREE
5215 344 : && operand_equal_p (OMP_CLAUSE_SIZE (c1),
5216 344 : OMP_CLAUSE_SIZE (c2))))
5217 82310 : && ((OMP_CLAUSE_ITERATORS (c1) == NULL_TREE
5218 523 : && OMP_CLAUSE_ITERATORS (c2) == NULL_TREE)
5219 0 : || operand_equal_p (OMP_CLAUSE_ITERATORS (c1),
5220 99773 : OMP_CLAUSE_ITERATORS (c2))));
5221 523 : if (maybe_dup_found)
5222 : {
5223 523 : if (first)
5224 : {
5225 234 : if (OMP_CLAUSE_MAP_KIND (c1) == OMP_CLAUSE_MAP_KIND (c2))
5226 : {
5227 : remove = true;
5228 : break;
5229 : }
5230 126 : else if ((OMP_CLAUSE_MAP_KIND (c1) & ~GOMP_MAP_TOFROM)
5231 126 : == (OMP_CLAUSE_MAP_KIND (c2) & ~GOMP_MAP_TOFROM))
5232 : {
5233 93 : OMP_CLAUSE_SET_MAP_KIND (c2,
5234 : (OMP_CLAUSE_MAP_KIND (c1)
5235 : | OMP_CLAUSE_MAP_KIND (c2)));
5236 93 : remove = true;
5237 93 : break;
5238 : }
5239 : }
5240 : /* When called from the gimplifier, remove duplicate map clauses
5241 : with identical kind only when the bits above
5242 : GOMP_MAP_FLAG_SPECIAL_2 are unset - as clauses with those flags
5243 : set may need to be present multiple times. */
5244 289 : else if (OMP_CLAUSE_MAP_KIND (c1) == OMP_CLAUSE_MAP_KIND (c2)
5245 289 : && (OMP_CLAUSE_MAP_KIND (c1) & ~0b11111) == 0)
5246 : {
5247 : remove = true;
5248 : break;
5249 : }
5250 : }
5251 : }
5252 24343 : if (remove)
5253 : remove = false;
5254 : else
5255 : {
5256 24131 : *outlist_p = c1;
5257 24131 : outlist_p = &OMP_CLAUSE_CHAIN (*outlist_p);
5258 : }
5259 : }
5260 70170 : *outlist_p = c1;
5261 70170 : return outlist;
5262 : }
|
