]> git.ipfire.org Git - thirdparty/gcc.git/blob - gcc/omp-general.c
arm: Factorize several occurrences of the same code into reg_needs_saving_p
[thirdparty/gcc.git] / gcc / omp-general.c
1 /* General types and functions that are uselful for processing of OpenMP,
2 OpenACC and similar directivers at various stages of compilation.
3
4 Copyright (C) 2005-2020 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 /* Find an OMP clause of type KIND within CLAUSES. */
23
24 #include "config.h"
25 #include "system.h"
26 #include "coretypes.h"
27 #include "backend.h"
28 #include "target.h"
29 #include "tree.h"
30 #include "gimple.h"
31 #include "ssa.h"
32 #include "diagnostic-core.h"
33 #include "fold-const.h"
34 #include "langhooks.h"
35 #include "omp-general.h"
36 #include "stringpool.h"
37 #include "attribs.h"
38 #include "gimplify.h"
39 #include "cgraph.h"
40 #include "alloc-pool.h"
41 #include "symbol-summary.h"
42 #include "hsa-common.h"
43 #include "tree-pass.h"
44 #include "omp-device-properties.h"
45
46 enum omp_requires omp_requires_mask;
47
48 tree
49 omp_find_clause (tree clauses, enum omp_clause_code kind)
50 {
51 for (; clauses ; clauses = OMP_CLAUSE_CHAIN (clauses))
52 if (OMP_CLAUSE_CODE (clauses) == kind)
53 return clauses;
54
55 return NULL_TREE;
56 }
57
58 /* True if OpenMP should regard this DECL as being a scalar which has Fortran's
59 allocatable or pointer attribute. */
60 bool
61 omp_is_allocatable_or_ptr (tree decl)
62 {
63 return lang_hooks.decls.omp_is_allocatable_or_ptr (decl);
64 }
65
66 /* Check whether this DECL belongs to a Fortran optional argument.
67 With 'for_present_check' set to false, decls which are optional parameters
68 themselve are returned as tree - or a NULL_TREE otherwise. Those decls are
69 always pointers. With 'for_present_check' set to true, the decl for checking
70 whether an argument is present is returned; for arguments with value
71 attribute this is the hidden argument and of BOOLEAN_TYPE. If the decl is
72 unrelated to optional arguments, NULL_TREE is returned. */
73
74 tree
75 omp_check_optional_argument (tree decl, bool for_present_check)
76 {
77 return lang_hooks.decls.omp_check_optional_argument (decl, for_present_check);
78 }
79
80 /* Return true if DECL is a reference type. */
81
82 bool
83 omp_is_reference (tree decl)
84 {
85 return lang_hooks.decls.omp_privatize_by_reference (decl);
86 }
87
88 /* Adjust *COND_CODE and *N2 so that the former is either LT_EXPR or GT_EXPR,
89 given that V is the loop index variable and STEP is loop step. */
90
91 void
92 omp_adjust_for_condition (location_t loc, enum tree_code *cond_code, tree *n2,
93 tree v, tree step)
94 {
95 switch (*cond_code)
96 {
97 case LT_EXPR:
98 case GT_EXPR:
99 break;
100
101 case NE_EXPR:
102 gcc_assert (TREE_CODE (step) == INTEGER_CST);
103 if (TREE_CODE (TREE_TYPE (v)) == INTEGER_TYPE)
104 {
105 if (integer_onep (step))
106 *cond_code = LT_EXPR;
107 else
108 {
109 gcc_assert (integer_minus_onep (step));
110 *cond_code = GT_EXPR;
111 }
112 }
113 else
114 {
115 tree unit = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (v)));
116 gcc_assert (TREE_CODE (unit) == INTEGER_CST);
117 if (tree_int_cst_equal (unit, step))
118 *cond_code = LT_EXPR;
119 else
120 {
121 gcc_assert (wi::neg (wi::to_widest (unit))
122 == wi::to_widest (step));
123 *cond_code = GT_EXPR;
124 }
125 }
126
127 break;
128
129 case LE_EXPR:
130 if (POINTER_TYPE_P (TREE_TYPE (*n2)))
131 *n2 = fold_build_pointer_plus_hwi_loc (loc, *n2, 1);
132 else
133 *n2 = fold_build2_loc (loc, PLUS_EXPR, TREE_TYPE (*n2), *n2,
134 build_int_cst (TREE_TYPE (*n2), 1));
135 *cond_code = LT_EXPR;
136 break;
137 case GE_EXPR:
138 if (POINTER_TYPE_P (TREE_TYPE (*n2)))
139 *n2 = fold_build_pointer_plus_hwi_loc (loc, *n2, -1);
140 else
141 *n2 = fold_build2_loc (loc, MINUS_EXPR, TREE_TYPE (*n2), *n2,
142 build_int_cst (TREE_TYPE (*n2), 1));
143 *cond_code = GT_EXPR;
144 break;
145 default:
146 gcc_unreachable ();
147 }
148 }
149
150 /* Return the looping step from INCR, extracted from the step of a gimple omp
151 for statement. */
152
153 tree
154 omp_get_for_step_from_incr (location_t loc, tree incr)
155 {
156 tree step;
157 switch (TREE_CODE (incr))
158 {
159 case PLUS_EXPR:
160 step = TREE_OPERAND (incr, 1);
161 break;
162 case POINTER_PLUS_EXPR:
163 step = fold_convert (ssizetype, TREE_OPERAND (incr, 1));
164 break;
165 case MINUS_EXPR:
166 step = TREE_OPERAND (incr, 1);
167 step = fold_build1_loc (loc, NEGATE_EXPR, TREE_TYPE (step), step);
168 break;
169 default:
170 gcc_unreachable ();
171 }
172 return step;
173 }
174
175 /* Extract the header elements of parallel loop FOR_STMT and store
176 them into *FD. */
177
178 void
179 omp_extract_for_data (gomp_for *for_stmt, struct omp_for_data *fd,
180 struct omp_for_data_loop *loops)
181 {
182 tree t, var, *collapse_iter, *collapse_count;
183 tree count = NULL_TREE, iter_type = long_integer_type_node;
184 struct omp_for_data_loop *loop;
185 int i;
186 struct omp_for_data_loop dummy_loop;
187 location_t loc = gimple_location (for_stmt);
188 bool simd = gimple_omp_for_kind (for_stmt) == GF_OMP_FOR_KIND_SIMD;
189 bool distribute = gimple_omp_for_kind (for_stmt)
190 == GF_OMP_FOR_KIND_DISTRIBUTE;
191 bool taskloop = gimple_omp_for_kind (for_stmt)
192 == GF_OMP_FOR_KIND_TASKLOOP;
193 tree iterv, countv;
194
195 fd->for_stmt = for_stmt;
196 fd->pre = NULL;
197 fd->have_nowait = distribute || simd;
198 fd->have_ordered = false;
199 fd->have_reductemp = false;
200 fd->have_pointer_condtemp = false;
201 fd->have_scantemp = false;
202 fd->have_nonctrl_scantemp = false;
203 fd->lastprivate_conditional = 0;
204 fd->tiling = NULL_TREE;
205 fd->collapse = 1;
206 fd->ordered = 0;
207 fd->sched_kind = OMP_CLAUSE_SCHEDULE_STATIC;
208 fd->sched_modifiers = 0;
209 fd->chunk_size = NULL_TREE;
210 fd->simd_schedule = false;
211 collapse_iter = NULL;
212 collapse_count = NULL;
213
214 for (t = gimple_omp_for_clauses (for_stmt); t ; t = OMP_CLAUSE_CHAIN (t))
215 switch (OMP_CLAUSE_CODE (t))
216 {
217 case OMP_CLAUSE_NOWAIT:
218 fd->have_nowait = true;
219 break;
220 case OMP_CLAUSE_ORDERED:
221 fd->have_ordered = true;
222 if (OMP_CLAUSE_ORDERED_EXPR (t))
223 fd->ordered = tree_to_shwi (OMP_CLAUSE_ORDERED_EXPR (t));
224 break;
225 case OMP_CLAUSE_SCHEDULE:
226 gcc_assert (!distribute && !taskloop);
227 fd->sched_kind
228 = (enum omp_clause_schedule_kind)
229 (OMP_CLAUSE_SCHEDULE_KIND (t) & OMP_CLAUSE_SCHEDULE_MASK);
230 fd->sched_modifiers = (OMP_CLAUSE_SCHEDULE_KIND (t)
231 & ~OMP_CLAUSE_SCHEDULE_MASK);
232 fd->chunk_size = OMP_CLAUSE_SCHEDULE_CHUNK_EXPR (t);
233 fd->simd_schedule = OMP_CLAUSE_SCHEDULE_SIMD (t);
234 break;
235 case OMP_CLAUSE_DIST_SCHEDULE:
236 gcc_assert (distribute);
237 fd->chunk_size = OMP_CLAUSE_DIST_SCHEDULE_CHUNK_EXPR (t);
238 break;
239 case OMP_CLAUSE_COLLAPSE:
240 fd->collapse = tree_to_shwi (OMP_CLAUSE_COLLAPSE_EXPR (t));
241 if (fd->collapse > 1)
242 {
243 collapse_iter = &OMP_CLAUSE_COLLAPSE_ITERVAR (t);
244 collapse_count = &OMP_CLAUSE_COLLAPSE_COUNT (t);
245 }
246 break;
247 case OMP_CLAUSE_TILE:
248 fd->tiling = OMP_CLAUSE_TILE_LIST (t);
249 fd->collapse = list_length (fd->tiling);
250 gcc_assert (fd->collapse);
251 collapse_iter = &OMP_CLAUSE_TILE_ITERVAR (t);
252 collapse_count = &OMP_CLAUSE_TILE_COUNT (t);
253 break;
254 case OMP_CLAUSE__REDUCTEMP_:
255 fd->have_reductemp = true;
256 break;
257 case OMP_CLAUSE_LASTPRIVATE:
258 if (OMP_CLAUSE_LASTPRIVATE_CONDITIONAL (t))
259 fd->lastprivate_conditional++;
260 break;
261 case OMP_CLAUSE__CONDTEMP_:
262 if (POINTER_TYPE_P (TREE_TYPE (OMP_CLAUSE_DECL (t))))
263 fd->have_pointer_condtemp = true;
264 break;
265 case OMP_CLAUSE__SCANTEMP_:
266 fd->have_scantemp = true;
267 if (!OMP_CLAUSE__SCANTEMP__ALLOC (t)
268 && !OMP_CLAUSE__SCANTEMP__CONTROL (t))
269 fd->have_nonctrl_scantemp = true;
270 break;
271 default:
272 break;
273 }
274
275 if (fd->collapse > 1 || fd->tiling)
276 fd->loops = loops;
277 else
278 fd->loops = &fd->loop;
279
280 if (fd->ordered && fd->collapse == 1 && loops != NULL)
281 {
282 fd->loops = loops;
283 iterv = NULL_TREE;
284 countv = NULL_TREE;
285 collapse_iter = &iterv;
286 collapse_count = &countv;
287 }
288
289 /* FIXME: for now map schedule(auto) to schedule(static).
290 There should be analysis to determine whether all iterations
291 are approximately the same amount of work (then schedule(static)
292 is best) or if it varies (then schedule(dynamic,N) is better). */
293 if (fd->sched_kind == OMP_CLAUSE_SCHEDULE_AUTO)
294 {
295 fd->sched_kind = OMP_CLAUSE_SCHEDULE_STATIC;
296 gcc_assert (fd->chunk_size == NULL);
297 }
298 gcc_assert ((fd->collapse == 1 && !fd->tiling) || collapse_iter != NULL);
299 if (taskloop)
300 fd->sched_kind = OMP_CLAUSE_SCHEDULE_RUNTIME;
301 if (fd->sched_kind == OMP_CLAUSE_SCHEDULE_RUNTIME)
302 gcc_assert (fd->chunk_size == NULL);
303 else if (fd->chunk_size == NULL)
304 {
305 /* We only need to compute a default chunk size for ordered
306 static loops and dynamic loops. */
307 if (fd->sched_kind != OMP_CLAUSE_SCHEDULE_STATIC
308 || fd->have_ordered)
309 fd->chunk_size = (fd->sched_kind == OMP_CLAUSE_SCHEDULE_STATIC)
310 ? integer_zero_node : integer_one_node;
311 }
312
313 int cnt = fd->ordered ? fd->ordered : fd->collapse;
314 for (i = 0; i < cnt; i++)
315 {
316 if (i == 0
317 && fd->collapse == 1
318 && !fd->tiling
319 && (fd->ordered == 0 || loops == NULL))
320 loop = &fd->loop;
321 else if (loops != NULL)
322 loop = loops + i;
323 else
324 loop = &dummy_loop;
325
326 loop->v = gimple_omp_for_index (for_stmt, i);
327 gcc_assert (SSA_VAR_P (loop->v));
328 gcc_assert (TREE_CODE (TREE_TYPE (loop->v)) == INTEGER_TYPE
329 || TREE_CODE (TREE_TYPE (loop->v)) == POINTER_TYPE);
330 var = TREE_CODE (loop->v) == SSA_NAME ? SSA_NAME_VAR (loop->v) : loop->v;
331 loop->n1 = gimple_omp_for_initial (for_stmt, i);
332
333 loop->cond_code = gimple_omp_for_cond (for_stmt, i);
334 loop->n2 = gimple_omp_for_final (for_stmt, i);
335 gcc_assert (loop->cond_code != NE_EXPR
336 || (gimple_omp_for_kind (for_stmt)
337 != GF_OMP_FOR_KIND_OACC_LOOP));
338
339 t = gimple_omp_for_incr (for_stmt, i);
340 gcc_assert (TREE_OPERAND (t, 0) == var);
341 loop->step = omp_get_for_step_from_incr (loc, t);
342
343 omp_adjust_for_condition (loc, &loop->cond_code, &loop->n2, loop->v,
344 loop->step);
345
346 if (simd
347 || (fd->sched_kind == OMP_CLAUSE_SCHEDULE_STATIC
348 && !fd->have_ordered))
349 {
350 if (fd->collapse == 1 && !fd->tiling)
351 iter_type = TREE_TYPE (loop->v);
352 else if (i == 0
353 || TYPE_PRECISION (iter_type)
354 < TYPE_PRECISION (TREE_TYPE (loop->v)))
355 iter_type
356 = build_nonstandard_integer_type
357 (TYPE_PRECISION (TREE_TYPE (loop->v)), 1);
358 }
359 else if (iter_type != long_long_unsigned_type_node)
360 {
361 if (POINTER_TYPE_P (TREE_TYPE (loop->v)))
362 iter_type = long_long_unsigned_type_node;
363 else if (TYPE_UNSIGNED (TREE_TYPE (loop->v))
364 && TYPE_PRECISION (TREE_TYPE (loop->v))
365 >= TYPE_PRECISION (iter_type))
366 {
367 tree n;
368
369 if (loop->cond_code == LT_EXPR)
370 n = fold_build2_loc (loc, PLUS_EXPR, TREE_TYPE (loop->v),
371 loop->n2, loop->step);
372 else
373 n = loop->n1;
374 if (TREE_CODE (n) != INTEGER_CST
375 || tree_int_cst_lt (TYPE_MAX_VALUE (iter_type), n))
376 iter_type = long_long_unsigned_type_node;
377 }
378 else if (TYPE_PRECISION (TREE_TYPE (loop->v))
379 > TYPE_PRECISION (iter_type))
380 {
381 tree n1, n2;
382
383 if (loop->cond_code == LT_EXPR)
384 {
385 n1 = loop->n1;
386 n2 = fold_build2_loc (loc, PLUS_EXPR, TREE_TYPE (loop->v),
387 loop->n2, loop->step);
388 }
389 else
390 {
391 n1 = fold_build2_loc (loc, MINUS_EXPR, TREE_TYPE (loop->v),
392 loop->n2, loop->step);
393 n2 = loop->n1;
394 }
395 if (TREE_CODE (n1) != INTEGER_CST
396 || TREE_CODE (n2) != INTEGER_CST
397 || !tree_int_cst_lt (TYPE_MIN_VALUE (iter_type), n1)
398 || !tree_int_cst_lt (n2, TYPE_MAX_VALUE (iter_type)))
399 iter_type = long_long_unsigned_type_node;
400 }
401 }
402
403 if (i >= fd->collapse)
404 continue;
405
406 if (collapse_count && *collapse_count == NULL)
407 {
408 t = fold_binary (loop->cond_code, boolean_type_node,
409 fold_convert (TREE_TYPE (loop->v), loop->n1),
410 fold_convert (TREE_TYPE (loop->v), loop->n2));
411 if (t && integer_zerop (t))
412 count = build_zero_cst (long_long_unsigned_type_node);
413 else if ((i == 0 || count != NULL_TREE)
414 && TREE_CODE (TREE_TYPE (loop->v)) == INTEGER_TYPE
415 && TREE_CONSTANT (loop->n1)
416 && TREE_CONSTANT (loop->n2)
417 && TREE_CODE (loop->step) == INTEGER_CST)
418 {
419 tree itype = TREE_TYPE (loop->v);
420
421 if (POINTER_TYPE_P (itype))
422 itype = signed_type_for (itype);
423 t = build_int_cst (itype, (loop->cond_code == LT_EXPR ? -1 : 1));
424 t = fold_build2_loc (loc, PLUS_EXPR, itype,
425 fold_convert_loc (loc, itype, loop->step),
426 t);
427 t = fold_build2_loc (loc, PLUS_EXPR, itype, t,
428 fold_convert_loc (loc, itype, loop->n2));
429 t = fold_build2_loc (loc, MINUS_EXPR, itype, t,
430 fold_convert_loc (loc, itype, loop->n1));
431 if (TYPE_UNSIGNED (itype) && loop->cond_code == GT_EXPR)
432 {
433 tree step = fold_convert_loc (loc, itype, loop->step);
434 t = fold_build2_loc (loc, TRUNC_DIV_EXPR, itype,
435 fold_build1_loc (loc, NEGATE_EXPR,
436 itype, t),
437 fold_build1_loc (loc, NEGATE_EXPR,
438 itype, step));
439 }
440 else
441 t = fold_build2_loc (loc, TRUNC_DIV_EXPR, itype, t,
442 fold_convert_loc (loc, itype,
443 loop->step));
444 t = fold_convert_loc (loc, long_long_unsigned_type_node, t);
445 if (count != NULL_TREE)
446 count = fold_build2_loc (loc, MULT_EXPR,
447 long_long_unsigned_type_node,
448 count, t);
449 else
450 count = t;
451 if (TREE_CODE (count) != INTEGER_CST)
452 count = NULL_TREE;
453 }
454 else if (count && !integer_zerop (count))
455 count = NULL_TREE;
456 }
457 }
458
459 if (count
460 && !simd
461 && (fd->sched_kind != OMP_CLAUSE_SCHEDULE_STATIC
462 || fd->have_ordered))
463 {
464 if (!tree_int_cst_lt (count, TYPE_MAX_VALUE (long_integer_type_node)))
465 iter_type = long_long_unsigned_type_node;
466 else
467 iter_type = long_integer_type_node;
468 }
469 else if (collapse_iter && *collapse_iter != NULL)
470 iter_type = TREE_TYPE (*collapse_iter);
471 fd->iter_type = iter_type;
472 if (collapse_iter && *collapse_iter == NULL)
473 *collapse_iter = create_tmp_var (iter_type, ".iter");
474 if (collapse_count && *collapse_count == NULL)
475 {
476 if (count)
477 *collapse_count = fold_convert_loc (loc, iter_type, count);
478 else
479 *collapse_count = create_tmp_var (iter_type, ".count");
480 }
481
482 if (fd->collapse > 1 || fd->tiling || (fd->ordered && loops))
483 {
484 fd->loop.v = *collapse_iter;
485 fd->loop.n1 = build_int_cst (TREE_TYPE (fd->loop.v), 0);
486 fd->loop.n2 = *collapse_count;
487 fd->loop.step = build_int_cst (TREE_TYPE (fd->loop.v), 1);
488 fd->loop.cond_code = LT_EXPR;
489 }
490 else if (loops)
491 loops[0] = fd->loop;
492 }
493
494 /* Build a call to GOMP_barrier. */
495
496 gimple *
497 omp_build_barrier (tree lhs)
498 {
499 tree fndecl = builtin_decl_explicit (lhs ? BUILT_IN_GOMP_BARRIER_CANCEL
500 : BUILT_IN_GOMP_BARRIER);
501 gcall *g = gimple_build_call (fndecl, 0);
502 if (lhs)
503 gimple_call_set_lhs (g, lhs);
504 return g;
505 }
506
507 /* Return maximum possible vectorization factor for the target. */
508
509 poly_uint64
510 omp_max_vf (void)
511 {
512 if (!optimize
513 || optimize_debug
514 || !flag_tree_loop_optimize
515 || (!flag_tree_loop_vectorize
516 && global_options_set.x_flag_tree_loop_vectorize))
517 return 1;
518
519 auto_vector_modes modes;
520 targetm.vectorize.autovectorize_vector_modes (&modes, true);
521 if (!modes.is_empty ())
522 {
523 poly_uint64 vf = 0;
524 for (unsigned int i = 0; i < modes.length (); ++i)
525 /* The returned modes use the smallest element size (and thus
526 the largest nunits) for the vectorization approach that they
527 represent. */
528 vf = ordered_max (vf, GET_MODE_NUNITS (modes[i]));
529 return vf;
530 }
531
532 machine_mode vqimode = targetm.vectorize.preferred_simd_mode (QImode);
533 if (GET_MODE_CLASS (vqimode) == MODE_VECTOR_INT)
534 return GET_MODE_NUNITS (vqimode);
535
536 return 1;
537 }
538
539 /* Return maximum SIMT width if offloading may target SIMT hardware. */
540
541 int
542 omp_max_simt_vf (void)
543 {
544 if (!optimize)
545 return 0;
546 if (ENABLE_OFFLOADING)
547 for (const char *c = getenv ("OFFLOAD_TARGET_NAMES"); c;)
548 {
549 if (!strncmp (c, "nvptx", strlen ("nvptx")))
550 return 32;
551 else if ((c = strchr (c, ':')))
552 c++;
553 }
554 return 0;
555 }
556
557 /* Store the construct selectors as tree codes from last to first,
558 return their number. */
559
560 int
561 omp_constructor_traits_to_codes (tree ctx, enum tree_code *constructs)
562 {
563 int nconstructs = list_length (ctx);
564 int i = nconstructs - 1;
565 for (tree t2 = ctx; t2; t2 = TREE_CHAIN (t2), i--)
566 {
567 const char *sel = IDENTIFIER_POINTER (TREE_PURPOSE (t2));
568 if (!strcmp (sel, "target"))
569 constructs[i] = OMP_TARGET;
570 else if (!strcmp (sel, "teams"))
571 constructs[i] = OMP_TEAMS;
572 else if (!strcmp (sel, "parallel"))
573 constructs[i] = OMP_PARALLEL;
574 else if (!strcmp (sel, "for") || !strcmp (sel, "do"))
575 constructs[i] = OMP_FOR;
576 else if (!strcmp (sel, "simd"))
577 constructs[i] = OMP_SIMD;
578 else
579 gcc_unreachable ();
580 }
581 gcc_assert (i == -1);
582 return nconstructs;
583 }
584
585 /* Return true if PROP is possibly present in one of the offloading target's
586 OpenMP contexts. The format of PROPS string is always offloading target's
587 name terminated by '\0', followed by properties for that offloading
588 target separated by '\0' and terminated by another '\0'. The strings
589 are created from omp-device-properties installed files of all configured
590 offloading targets. */
591
592 static bool
593 omp_offload_device_kind_arch_isa (const char *props, const char *prop)
594 {
595 const char *names = getenv ("OFFLOAD_TARGET_NAMES");
596 if (names == NULL || *names == '\0')
597 return false;
598 while (*props != '\0')
599 {
600 size_t name_len = strlen (props);
601 bool matches = false;
602 for (const char *c = names; c; )
603 {
604 if (strncmp (props, c, name_len) == 0
605 && (c[name_len] == '\0'
606 || c[name_len] == ':'
607 || c[name_len] == '='))
608 {
609 matches = true;
610 break;
611 }
612 else if ((c = strchr (c, ':')))
613 c++;
614 }
615 props = props + name_len + 1;
616 while (*props != '\0')
617 {
618 if (matches && strcmp (props, prop) == 0)
619 return true;
620 props = strchr (props, '\0') + 1;
621 }
622 props++;
623 }
624 return false;
625 }
626
627 /* Return true if the current code location is or might be offloaded.
628 Return true in declare target functions, or when nested in a target
629 region or when unsure, return false otherwise. */
630
631 static bool
632 omp_maybe_offloaded (void)
633 {
634 if (!hsa_gen_requested_p ())
635 {
636 if (!ENABLE_OFFLOADING)
637 return false;
638 const char *names = getenv ("OFFLOAD_TARGET_NAMES");
639 if (names == NULL || *names == '\0')
640 return false;
641 }
642 if (symtab->state == PARSING)
643 /* Maybe. */
644 return true;
645 if (cfun && cfun->after_inlining)
646 return false;
647 if (current_function_decl
648 && lookup_attribute ("omp declare target",
649 DECL_ATTRIBUTES (current_function_decl)))
650 return true;
651 if (cfun && (cfun->curr_properties & PROP_gimple_any) == 0)
652 {
653 enum tree_code construct = OMP_TARGET;
654 if (omp_construct_selector_matches (&construct, 1, NULL))
655 return true;
656 }
657 return false;
658 }
659
660 /* Return a name from PROP, a property in selectors accepting
661 name lists. */
662
663 static const char *
664 omp_context_name_list_prop (tree prop)
665 {
666 if (TREE_PURPOSE (prop))
667 return IDENTIFIER_POINTER (TREE_PURPOSE (prop));
668 else
669 {
670 const char *ret = TREE_STRING_POINTER (TREE_VALUE (prop));
671 if ((size_t) TREE_STRING_LENGTH (TREE_VALUE (prop)) == strlen (ret) + 1)
672 return ret;
673 return NULL;
674 }
675 }
676
677 /* Return 1 if context selector matches the current OpenMP context, 0
678 if it does not and -1 if it is unknown and need to be determined later.
679 Some properties can be checked right away during parsing (this routine),
680 others need to wait until the whole TU is parsed, others need to wait until
681 IPA, others until vectorization. */
682
683 int
684 omp_context_selector_matches (tree ctx)
685 {
686 int ret = 1;
687 for (tree t1 = ctx; t1; t1 = TREE_CHAIN (t1))
688 {
689 char set = IDENTIFIER_POINTER (TREE_PURPOSE (t1))[0];
690 if (set == 'c')
691 {
692 /* For now, ignore the construct set. While something can be
693 determined already during parsing, we don't know until end of TU
694 whether additional constructs aren't added through declare variant
695 unless "omp declare variant variant" attribute exists already
696 (so in most of the cases), and we'd need to maintain set of
697 surrounding OpenMP constructs, which is better handled during
698 gimplification. */
699 if (symtab->state == PARSING)
700 {
701 ret = -1;
702 continue;
703 }
704
705 enum tree_code constructs[5];
706 int nconstructs
707 = omp_constructor_traits_to_codes (TREE_VALUE (t1), constructs);
708
709 if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
710 {
711 if (!cfun->after_inlining)
712 {
713 ret = -1;
714 continue;
715 }
716 int i;
717 for (i = 0; i < nconstructs; ++i)
718 if (constructs[i] == OMP_SIMD)
719 break;
720 if (i < nconstructs)
721 {
722 ret = -1;
723 continue;
724 }
725 /* If there is no simd, assume it is ok after IPA,
726 constructs should have been checked before. */
727 continue;
728 }
729
730 int r = omp_construct_selector_matches (constructs, nconstructs,
731 NULL);
732 if (r == 0)
733 return 0;
734 if (r == -1)
735 ret = -1;
736 continue;
737 }
738 for (tree t2 = TREE_VALUE (t1); t2; t2 = TREE_CHAIN (t2))
739 {
740 const char *sel = IDENTIFIER_POINTER (TREE_PURPOSE (t2));
741 switch (*sel)
742 {
743 case 'v':
744 if (set == 'i' && !strcmp (sel, "vendor"))
745 for (tree t3 = TREE_VALUE (t2); t3; t3 = TREE_CHAIN (t3))
746 {
747 const char *prop = omp_context_name_list_prop (t3);
748 if (prop == NULL)
749 return 0;
750 if ((!strcmp (prop, " score") && TREE_PURPOSE (t3))
751 || !strcmp (prop, "gnu"))
752 continue;
753 return 0;
754 }
755 break;
756 case 'e':
757 if (set == 'i' && !strcmp (sel, "extension"))
758 /* We don't support any extensions right now. */
759 return 0;
760 break;
761 case 'a':
762 if (set == 'i' && !strcmp (sel, "atomic_default_mem_order"))
763 {
764 if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
765 break;
766
767 enum omp_memory_order omo
768 = ((enum omp_memory_order)
769 (omp_requires_mask
770 & OMP_REQUIRES_ATOMIC_DEFAULT_MEM_ORDER));
771 if (omo == OMP_MEMORY_ORDER_UNSPECIFIED)
772 {
773 /* We don't know yet, until end of TU. */
774 if (symtab->state == PARSING)
775 {
776 ret = -1;
777 break;
778 }
779 else
780 omo = OMP_MEMORY_ORDER_RELAXED;
781 }
782 tree t3 = TREE_VALUE (t2);
783 const char *prop = IDENTIFIER_POINTER (TREE_PURPOSE (t3));
784 if (!strcmp (prop, " score"))
785 {
786 t3 = TREE_CHAIN (t3);
787 prop = IDENTIFIER_POINTER (TREE_PURPOSE (t3));
788 }
789 if (!strcmp (prop, "relaxed")
790 && omo != OMP_MEMORY_ORDER_RELAXED)
791 return 0;
792 else if (!strcmp (prop, "seq_cst")
793 && omo != OMP_MEMORY_ORDER_SEQ_CST)
794 return 0;
795 else if (!strcmp (prop, "acq_rel")
796 && omo != OMP_MEMORY_ORDER_ACQ_REL)
797 return 0;
798 }
799 if (set == 'd' && !strcmp (sel, "arch"))
800 for (tree t3 = TREE_VALUE (t2); t3; t3 = TREE_CHAIN (t3))
801 {
802 const char *arch = omp_context_name_list_prop (t3);
803 if (arch == NULL)
804 return 0;
805 int r = 0;
806 if (targetm.omp.device_kind_arch_isa != NULL)
807 r = targetm.omp.device_kind_arch_isa (omp_device_arch,
808 arch);
809 if (r == 0 || (r == -1 && symtab->state != PARSING))
810 {
811 /* If we are or might be in a target region or
812 declare target function, need to take into account
813 also offloading values. */
814 if (!omp_maybe_offloaded ())
815 return 0;
816 if (strcmp (arch, "hsa") == 0
817 && hsa_gen_requested_p ())
818 {
819 ret = -1;
820 continue;
821 }
822 if (ENABLE_OFFLOADING)
823 {
824 const char *arches = omp_offload_device_arch;
825 if (omp_offload_device_kind_arch_isa (arches,
826 arch))
827 {
828 ret = -1;
829 continue;
830 }
831 }
832 return 0;
833 }
834 else if (r == -1)
835 ret = -1;
836 /* If arch matches on the host, it still might not match
837 in the offloading region. */
838 else if (omp_maybe_offloaded ())
839 ret = -1;
840 }
841 break;
842 case 'u':
843 if (set == 'i' && !strcmp (sel, "unified_address"))
844 {
845 if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
846 break;
847
848 if ((omp_requires_mask & OMP_REQUIRES_UNIFIED_ADDRESS) == 0)
849 {
850 if (symtab->state == PARSING)
851 ret = -1;
852 else
853 return 0;
854 }
855 break;
856 }
857 if (set == 'i' && !strcmp (sel, "unified_shared_memory"))
858 {
859 if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
860 break;
861
862 if ((omp_requires_mask
863 & OMP_REQUIRES_UNIFIED_SHARED_MEMORY) == 0)
864 {
865 if (symtab->state == PARSING)
866 ret = -1;
867 else
868 return 0;
869 }
870 break;
871 }
872 break;
873 case 'd':
874 if (set == 'i' && !strcmp (sel, "dynamic_allocators"))
875 {
876 if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
877 break;
878
879 if ((omp_requires_mask
880 & OMP_REQUIRES_DYNAMIC_ALLOCATORS) == 0)
881 {
882 if (symtab->state == PARSING)
883 ret = -1;
884 else
885 return 0;
886 }
887 break;
888 }
889 break;
890 case 'r':
891 if (set == 'i' && !strcmp (sel, "reverse_offload"))
892 {
893 if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
894 break;
895
896 if ((omp_requires_mask & OMP_REQUIRES_REVERSE_OFFLOAD) == 0)
897 {
898 if (symtab->state == PARSING)
899 ret = -1;
900 else
901 return 0;
902 }
903 break;
904 }
905 break;
906 case 'k':
907 if (set == 'd' && !strcmp (sel, "kind"))
908 for (tree t3 = TREE_VALUE (t2); t3; t3 = TREE_CHAIN (t3))
909 {
910 const char *prop = omp_context_name_list_prop (t3);
911 if (prop == NULL)
912 return 0;
913 if (!strcmp (prop, "any"))
914 continue;
915 if (!strcmp (prop, "host"))
916 {
917 if (omp_maybe_offloaded ())
918 ret = -1;
919 continue;
920 }
921 if (!strcmp (prop, "nohost"))
922 {
923 if (omp_maybe_offloaded ())
924 ret = -1;
925 else
926 return 0;
927 continue;
928 }
929 int r = 0;
930 if (targetm.omp.device_kind_arch_isa != NULL)
931 r = targetm.omp.device_kind_arch_isa (omp_device_kind,
932 prop);
933 else
934 r = strcmp (prop, "cpu") == 0;
935 if (r == 0 || (r == -1 && symtab->state != PARSING))
936 {
937 /* If we are or might be in a target region or
938 declare target function, need to take into account
939 also offloading values. */
940 if (!omp_maybe_offloaded ())
941 return 0;
942 if (strcmp (prop, "gpu") == 0
943 && hsa_gen_requested_p ())
944 {
945 ret = -1;
946 continue;
947 }
948 if (ENABLE_OFFLOADING)
949 {
950 const char *kinds = omp_offload_device_kind;
951 if (omp_offload_device_kind_arch_isa (kinds, prop))
952 {
953 ret = -1;
954 continue;
955 }
956 }
957 return 0;
958 }
959 else if (r == -1)
960 ret = -1;
961 /* If kind matches on the host, it still might not match
962 in the offloading region. */
963 else if (omp_maybe_offloaded ())
964 ret = -1;
965 }
966 break;
967 case 'i':
968 if (set == 'd' && !strcmp (sel, "isa"))
969 for (tree t3 = TREE_VALUE (t2); t3; t3 = TREE_CHAIN (t3))
970 {
971 const char *isa = omp_context_name_list_prop (t3);
972 if (isa == NULL)
973 return 0;
974 int r = 0;
975 if (targetm.omp.device_kind_arch_isa != NULL)
976 r = targetm.omp.device_kind_arch_isa (omp_device_isa,
977 isa);
978 if (r == 0 || (r == -1 && symtab->state != PARSING))
979 {
980 /* If isa is valid on the target, but not in the
981 current function and current function has
982 #pragma omp declare simd on it, some simd clones
983 might have the isa added later on. */
984 if (r == -1
985 && targetm.simd_clone.compute_vecsize_and_simdlen
986 && (cfun == NULL || !cfun->after_inlining))
987 {
988 tree attrs
989 = DECL_ATTRIBUTES (current_function_decl);
990 if (lookup_attribute ("omp declare simd", attrs))
991 {
992 ret = -1;
993 continue;
994 }
995 }
996 /* If we are or might be in a target region or
997 declare target function, need to take into account
998 also offloading values. */
999 if (!omp_maybe_offloaded ())
1000 return 0;
1001 if (ENABLE_OFFLOADING)
1002 {
1003 const char *isas = omp_offload_device_isa;
1004 if (omp_offload_device_kind_arch_isa (isas, isa))
1005 {
1006 ret = -1;
1007 continue;
1008 }
1009 }
1010 return 0;
1011 }
1012 else if (r == -1)
1013 ret = -1;
1014 /* If isa matches on the host, it still might not match
1015 in the offloading region. */
1016 else if (omp_maybe_offloaded ())
1017 ret = -1;
1018 }
1019 break;
1020 case 'c':
1021 if (set == 'u' && !strcmp (sel, "condition"))
1022 for (tree t3 = TREE_VALUE (t2); t3; t3 = TREE_CHAIN (t3))
1023 if (TREE_PURPOSE (t3) == NULL_TREE)
1024 {
1025 if (integer_zerop (TREE_VALUE (t3)))
1026 return 0;
1027 if (integer_nonzerop (TREE_VALUE (t3)))
1028 break;
1029 ret = -1;
1030 }
1031 break;
1032 default:
1033 break;
1034 }
1035 }
1036 }
1037 return ret;
1038 }
1039
1040 /* Compare construct={simd} CLAUSES1 with CLAUSES2, return 0/-1/1/2 as
1041 in omp_context_selector_set_compare. */
1042
1043 static int
1044 omp_construct_simd_compare (tree clauses1, tree clauses2)
1045 {
1046 if (clauses1 == NULL_TREE)
1047 return clauses2 == NULL_TREE ? 0 : -1;
1048 if (clauses2 == NULL_TREE)
1049 return 1;
1050
1051 int r = 0;
1052 struct declare_variant_simd_data {
1053 bool inbranch, notinbranch;
1054 tree simdlen;
1055 auto_vec<tree,16> data_sharing;
1056 auto_vec<tree,16> aligned;
1057 declare_variant_simd_data ()
1058 : inbranch(false), notinbranch(false), simdlen(NULL_TREE) {}
1059 } data[2];
1060 unsigned int i;
1061 for (i = 0; i < 2; i++)
1062 for (tree c = i ? clauses2 : clauses1; c; c = OMP_CLAUSE_CHAIN (c))
1063 {
1064 vec<tree> *v;
1065 switch (OMP_CLAUSE_CODE (c))
1066 {
1067 case OMP_CLAUSE_INBRANCH:
1068 data[i].inbranch = true;
1069 continue;
1070 case OMP_CLAUSE_NOTINBRANCH:
1071 data[i].notinbranch = true;
1072 continue;
1073 case OMP_CLAUSE_SIMDLEN:
1074 data[i].simdlen = OMP_CLAUSE_SIMDLEN_EXPR (c);
1075 continue;
1076 case OMP_CLAUSE_UNIFORM:
1077 case OMP_CLAUSE_LINEAR:
1078 v = &data[i].data_sharing;
1079 break;
1080 case OMP_CLAUSE_ALIGNED:
1081 v = &data[i].aligned;
1082 break;
1083 default:
1084 gcc_unreachable ();
1085 }
1086 unsigned HOST_WIDE_INT argno = tree_to_uhwi (OMP_CLAUSE_DECL (c));
1087 if (argno >= v->length ())
1088 v->safe_grow_cleared (argno + 1);
1089 (*v)[argno] = c;
1090 }
1091 /* Here, r is used as a bitmask, 2 is set if CLAUSES1 has something
1092 CLAUSES2 doesn't, 1 is set if CLAUSES2 has something CLAUSES1
1093 doesn't. Thus, r == 3 implies return value 2, r == 1 implies
1094 -1, r == 2 implies 1 and r == 0 implies 0. */
1095 if (data[0].inbranch != data[1].inbranch)
1096 r |= data[0].inbranch ? 2 : 1;
1097 if (data[0].notinbranch != data[1].notinbranch)
1098 r |= data[0].notinbranch ? 2 : 1;
1099 if (!simple_cst_equal (data[0].simdlen, data[1].simdlen))
1100 {
1101 if (data[0].simdlen && data[1].simdlen)
1102 return 2;
1103 r |= data[0].simdlen ? 2 : 1;
1104 }
1105 if (data[0].data_sharing.length () < data[1].data_sharing.length ()
1106 || data[0].aligned.length () < data[1].aligned.length ())
1107 r |= 1;
1108 tree c1, c2;
1109 FOR_EACH_VEC_ELT (data[0].data_sharing, i, c1)
1110 {
1111 c2 = (i < data[1].data_sharing.length ()
1112 ? data[1].data_sharing[i] : NULL_TREE);
1113 if ((c1 == NULL_TREE) != (c2 == NULL_TREE))
1114 {
1115 r |= c1 != NULL_TREE ? 2 : 1;
1116 continue;
1117 }
1118 if (c1 == NULL_TREE)
1119 continue;
1120 if (OMP_CLAUSE_CODE (c1) != OMP_CLAUSE_CODE (c2))
1121 return 2;
1122 if (OMP_CLAUSE_CODE (c1) != OMP_CLAUSE_LINEAR)
1123 continue;
1124 if (OMP_CLAUSE_LINEAR_VARIABLE_STRIDE (c1)
1125 != OMP_CLAUSE_LINEAR_VARIABLE_STRIDE (c2))
1126 return 2;
1127 if (OMP_CLAUSE_LINEAR_KIND (c1) != OMP_CLAUSE_LINEAR_KIND (c2))
1128 return 2;
1129 if (!simple_cst_equal (OMP_CLAUSE_LINEAR_STEP (c1),
1130 OMP_CLAUSE_LINEAR_STEP (c2)))
1131 return 2;
1132 }
1133 FOR_EACH_VEC_ELT (data[0].aligned, i, c1)
1134 {
1135 c2 = i < data[1].aligned.length () ? data[1].aligned[i] : NULL_TREE;
1136 if ((c1 == NULL_TREE) != (c2 == NULL_TREE))
1137 {
1138 r |= c1 != NULL_TREE ? 2 : 1;
1139 continue;
1140 }
1141 if (c1 == NULL_TREE)
1142 continue;
1143 if (!simple_cst_equal (OMP_CLAUSE_ALIGNED_ALIGNMENT (c1),
1144 OMP_CLAUSE_ALIGNED_ALIGNMENT (c2)))
1145 return 2;
1146 }
1147 switch (r)
1148 {
1149 case 0: return 0;
1150 case 1: return -1;
1151 case 2: return 1;
1152 case 3: return 2;
1153 default: gcc_unreachable ();
1154 }
1155 }
1156
1157 /* Compare properties of selectors SEL from SET other than construct.
1158 Return 0/-1/1/2 as in omp_context_selector_set_compare.
1159 Unlike set names or selector names, properties can have duplicates. */
1160
1161 static int
1162 omp_context_selector_props_compare (const char *set, const char *sel,
1163 tree ctx1, tree ctx2)
1164 {
1165 int ret = 0;
1166 for (int pass = 0; pass < 2; pass++)
1167 for (tree t1 = pass ? ctx2 : ctx1; t1; t1 = TREE_CHAIN (t1))
1168 {
1169 tree t2;
1170 for (t2 = pass ? ctx1 : ctx2; t2; t2 = TREE_CHAIN (t2))
1171 if (TREE_PURPOSE (t1) == TREE_PURPOSE (t2))
1172 {
1173 if (TREE_PURPOSE (t1) == NULL_TREE)
1174 {
1175 if (set[0] == 'u' && strcmp (sel, "condition") == 0)
1176 {
1177 if (integer_zerop (TREE_VALUE (t1))
1178 != integer_zerop (TREE_VALUE (t2)))
1179 return 2;
1180 break;
1181 }
1182 if (simple_cst_equal (TREE_VALUE (t1), TREE_VALUE (t2)))
1183 break;
1184 }
1185 else if (strcmp (IDENTIFIER_POINTER (TREE_PURPOSE (t1)),
1186 " score") == 0)
1187 {
1188 if (!simple_cst_equal (TREE_VALUE (t1), TREE_VALUE (t2)))
1189 return 2;
1190 break;
1191 }
1192 else
1193 break;
1194 }
1195 else if (TREE_PURPOSE (t1)
1196 && TREE_PURPOSE (t2) == NULL_TREE
1197 && TREE_CODE (TREE_VALUE (t2)) == STRING_CST)
1198 {
1199 const char *p1 = omp_context_name_list_prop (t1);
1200 const char *p2 = omp_context_name_list_prop (t2);
1201 if (p2
1202 && strcmp (p1, p2) == 0
1203 && strcmp (p1, " score"))
1204 break;
1205 }
1206 else if (TREE_PURPOSE (t1) == NULL_TREE
1207 && TREE_PURPOSE (t2)
1208 && TREE_CODE (TREE_VALUE (t1)) == STRING_CST)
1209 {
1210 const char *p1 = omp_context_name_list_prop (t1);
1211 const char *p2 = omp_context_name_list_prop (t2);
1212 if (p1
1213 && strcmp (p1, p2) == 0
1214 && strcmp (p1, " score"))
1215 break;
1216 }
1217 if (t2 == NULL_TREE)
1218 {
1219 int r = pass ? -1 : 1;
1220 if (ret && ret != r)
1221 return 2;
1222 else if (pass)
1223 return r;
1224 else
1225 {
1226 ret = r;
1227 break;
1228 }
1229 }
1230 }
1231 return ret;
1232 }
1233
1234 /* Compare single context selector sets CTX1 and CTX2 with SET name.
1235 Return 0 if CTX1 is equal to CTX2,
1236 -1 if CTX1 is a strict subset of CTX2,
1237 1 if CTX2 is a strict subset of CTX1, or
1238 2 if neither context is a subset of another one. */
1239
1240 int
1241 omp_context_selector_set_compare (const char *set, tree ctx1, tree ctx2)
1242 {
1243 bool swapped = false;
1244 int ret = 0;
1245 int len1 = list_length (ctx1);
1246 int len2 = list_length (ctx2);
1247 int cnt = 0;
1248 if (len1 < len2)
1249 {
1250 swapped = true;
1251 std::swap (ctx1, ctx2);
1252 std::swap (len1, len2);
1253 }
1254 if (set[0] == 'c')
1255 {
1256 tree t1;
1257 tree t2 = ctx2;
1258 tree simd = get_identifier ("simd");
1259 /* Handle construct set specially. In this case the order
1260 of the selector matters too. */
1261 for (t1 = ctx1; t1; t1 = TREE_CHAIN (t1))
1262 if (TREE_PURPOSE (t1) == TREE_PURPOSE (t2))
1263 {
1264 int r = 0;
1265 if (TREE_PURPOSE (t1) == simd)
1266 r = omp_construct_simd_compare (TREE_VALUE (t1),
1267 TREE_VALUE (t2));
1268 if (r == 2 || (ret && r && (ret < 0) != (r < 0)))
1269 return 2;
1270 if (ret == 0)
1271 ret = r;
1272 t2 = TREE_CHAIN (t2);
1273 if (t2 == NULL_TREE)
1274 {
1275 t1 = TREE_CHAIN (t1);
1276 break;
1277 }
1278 }
1279 else if (ret < 0)
1280 return 2;
1281 else
1282 ret = 1;
1283 if (t2 != NULL_TREE)
1284 return 2;
1285 if (t1 != NULL_TREE)
1286 {
1287 if (ret < 0)
1288 return 2;
1289 ret = 1;
1290 }
1291 if (ret == 0)
1292 return 0;
1293 return swapped ? -ret : ret;
1294 }
1295 for (tree t1 = ctx1; t1; t1 = TREE_CHAIN (t1))
1296 {
1297 tree t2;
1298 for (t2 = ctx2; t2; t2 = TREE_CHAIN (t2))
1299 if (TREE_PURPOSE (t1) == TREE_PURPOSE (t2))
1300 {
1301 const char *sel = IDENTIFIER_POINTER (TREE_PURPOSE (t1));
1302 int r = omp_context_selector_props_compare (set, sel,
1303 TREE_VALUE (t1),
1304 TREE_VALUE (t2));
1305 if (r == 2 || (ret && r && (ret < 0) != (r < 0)))
1306 return 2;
1307 if (ret == 0)
1308 ret = r;
1309 cnt++;
1310 break;
1311 }
1312 if (t2 == NULL_TREE)
1313 {
1314 if (ret == -1)
1315 return 2;
1316 ret = 1;
1317 }
1318 }
1319 if (cnt < len2)
1320 return 2;
1321 if (ret == 0)
1322 return 0;
1323 return swapped ? -ret : ret;
1324 }
1325
1326 /* Compare whole context selector specification CTX1 and CTX2.
1327 Return 0 if CTX1 is equal to CTX2,
1328 -1 if CTX1 is a strict subset of CTX2,
1329 1 if CTX2 is a strict subset of CTX1, or
1330 2 if neither context is a subset of another one. */
1331
1332 static int
1333 omp_context_selector_compare (tree ctx1, tree ctx2)
1334 {
1335 bool swapped = false;
1336 int ret = 0;
1337 int len1 = list_length (ctx1);
1338 int len2 = list_length (ctx2);
1339 int cnt = 0;
1340 if (len1 < len2)
1341 {
1342 swapped = true;
1343 std::swap (ctx1, ctx2);
1344 std::swap (len1, len2);
1345 }
1346 for (tree t1 = ctx1; t1; t1 = TREE_CHAIN (t1))
1347 {
1348 tree t2;
1349 for (t2 = ctx2; t2; t2 = TREE_CHAIN (t2))
1350 if (TREE_PURPOSE (t1) == TREE_PURPOSE (t2))
1351 {
1352 const char *set = IDENTIFIER_POINTER (TREE_PURPOSE (t1));
1353 int r = omp_context_selector_set_compare (set, TREE_VALUE (t1),
1354 TREE_VALUE (t2));
1355 if (r == 2 || (ret && r && (ret < 0) != (r < 0)))
1356 return 2;
1357 if (ret == 0)
1358 ret = r;
1359 cnt++;
1360 break;
1361 }
1362 if (t2 == NULL_TREE)
1363 {
1364 if (ret == -1)
1365 return 2;
1366 ret = 1;
1367 }
1368 }
1369 if (cnt < len2)
1370 return 2;
1371 if (ret == 0)
1372 return 0;
1373 return swapped ? -ret : ret;
1374 }
1375
1376 /* From context selector CTX, return trait-selector with name SEL in
1377 trait-selector-set with name SET if any, or NULL_TREE if not found.
1378 If SEL is NULL, return the list of trait-selectors in SET. */
1379
1380 tree
1381 omp_get_context_selector (tree ctx, const char *set, const char *sel)
1382 {
1383 tree setid = get_identifier (set);
1384 tree selid = sel ? get_identifier (sel) : NULL_TREE;
1385 for (tree t1 = ctx; t1; t1 = TREE_CHAIN (t1))
1386 if (TREE_PURPOSE (t1) == setid)
1387 {
1388 if (sel == NULL)
1389 return TREE_VALUE (t1);
1390 for (tree t2 = TREE_VALUE (t1); t2; t2 = TREE_CHAIN (t2))
1391 if (TREE_PURPOSE (t2) == selid)
1392 return t2;
1393 }
1394 return NULL_TREE;
1395 }
1396
1397 /* Compute *SCORE for context selector CTX. Return true if the score
1398 would be different depending on whether it is a declare simd clone or
1399 not. DECLARE_SIMD should be true for the case when it would be
1400 a declare simd clone. */
1401
1402 static bool
1403 omp_context_compute_score (tree ctx, widest_int *score, bool declare_simd)
1404 {
1405 tree construct = omp_get_context_selector (ctx, "construct", NULL);
1406 bool has_kind = omp_get_context_selector (ctx, "device", "kind");
1407 bool has_arch = omp_get_context_selector (ctx, "device", "arch");
1408 bool has_isa = omp_get_context_selector (ctx, "device", "isa");
1409 bool ret = false;
1410 *score = 1;
1411 for (tree t1 = ctx; t1; t1 = TREE_CHAIN (t1))
1412 if (TREE_VALUE (t1) != construct)
1413 for (tree t2 = TREE_VALUE (t1); t2; t2 = TREE_CHAIN (t2))
1414 if (tree t3 = TREE_VALUE (t2))
1415 if (TREE_PURPOSE (t3)
1416 && strcmp (IDENTIFIER_POINTER (TREE_PURPOSE (t3)), " score") == 0
1417 && TREE_CODE (TREE_VALUE (t3)) == INTEGER_CST)
1418 *score += wi::to_widest (TREE_VALUE (t3));
1419 if (construct || has_kind || has_arch || has_isa)
1420 {
1421 int scores[12];
1422 enum tree_code constructs[5];
1423 int nconstructs = 0;
1424 if (construct)
1425 nconstructs = omp_constructor_traits_to_codes (construct, constructs);
1426 if (omp_construct_selector_matches (constructs, nconstructs, scores)
1427 == 2)
1428 ret = true;
1429 int b = declare_simd ? nconstructs + 1 : 0;
1430 if (scores[b + nconstructs] + 4U < score->get_precision ())
1431 {
1432 for (int n = 0; n < nconstructs; ++n)
1433 {
1434 if (scores[b + n] < 0)
1435 {
1436 *score = -1;
1437 return ret;
1438 }
1439 *score += wi::shifted_mask <widest_int> (scores[b + n], 1, false);
1440 }
1441 if (has_kind)
1442 *score += wi::shifted_mask <widest_int> (scores[b + nconstructs],
1443 1, false);
1444 if (has_arch)
1445 *score += wi::shifted_mask <widest_int> (scores[b + nconstructs] + 1,
1446 1, false);
1447 if (has_isa)
1448 *score += wi::shifted_mask <widest_int> (scores[b + nconstructs] + 2,
1449 1, false);
1450 }
1451 else /* FIXME: Implement this. */
1452 gcc_unreachable ();
1453 }
1454 return ret;
1455 }
1456
1457 /* Class describing a single variant. */
1458 struct GTY(()) omp_declare_variant_entry {
1459 /* NODE of the variant. */
1460 cgraph_node *variant;
1461 /* Score if not in declare simd clone. */
1462 widest_int score;
1463 /* Score if in declare simd clone. */
1464 widest_int score_in_declare_simd_clone;
1465 /* Context selector for the variant. */
1466 tree ctx;
1467 /* True if the context selector is known to match already. */
1468 bool matches;
1469 };
1470
1471 /* Class describing a function with variants. */
1472 struct GTY((for_user)) omp_declare_variant_base_entry {
1473 /* NODE of the base function. */
1474 cgraph_node *base;
1475 /* NODE of the artificial function created for the deferred variant
1476 resolution. */
1477 cgraph_node *node;
1478 /* Vector of the variants. */
1479 vec<omp_declare_variant_entry, va_gc> *variants;
1480 };
1481
1482 struct omp_declare_variant_hasher
1483 : ggc_ptr_hash<omp_declare_variant_base_entry> {
1484 static hashval_t hash (omp_declare_variant_base_entry *);
1485 static bool equal (omp_declare_variant_base_entry *,
1486 omp_declare_variant_base_entry *);
1487 };
1488
1489 hashval_t
1490 omp_declare_variant_hasher::hash (omp_declare_variant_base_entry *x)
1491 {
1492 inchash::hash hstate;
1493 hstate.add_int (DECL_UID (x->base->decl));
1494 hstate.add_int (x->variants->length ());
1495 omp_declare_variant_entry *variant;
1496 unsigned int i;
1497 FOR_EACH_VEC_SAFE_ELT (x->variants, i, variant)
1498 {
1499 hstate.add_int (DECL_UID (variant->variant->decl));
1500 hstate.add_wide_int (variant->score);
1501 hstate.add_wide_int (variant->score_in_declare_simd_clone);
1502 hstate.add_ptr (variant->ctx);
1503 hstate.add_int (variant->matches);
1504 }
1505 return hstate.end ();
1506 }
1507
1508 bool
1509 omp_declare_variant_hasher::equal (omp_declare_variant_base_entry *x,
1510 omp_declare_variant_base_entry *y)
1511 {
1512 if (x->base != y->base
1513 || x->variants->length () != y->variants->length ())
1514 return false;
1515 omp_declare_variant_entry *variant;
1516 unsigned int i;
1517 FOR_EACH_VEC_SAFE_ELT (x->variants, i, variant)
1518 if (variant->variant != (*y->variants)[i].variant
1519 || variant->score != (*y->variants)[i].score
1520 || (variant->score_in_declare_simd_clone
1521 != (*y->variants)[i].score_in_declare_simd_clone)
1522 || variant->ctx != (*y->variants)[i].ctx
1523 || variant->matches != (*y->variants)[i].matches)
1524 return false;
1525 return true;
1526 }
1527
1528 static GTY(()) hash_table<omp_declare_variant_hasher> *omp_declare_variants;
1529
1530 struct omp_declare_variant_alt_hasher
1531 : ggc_ptr_hash<omp_declare_variant_base_entry> {
1532 static hashval_t hash (omp_declare_variant_base_entry *);
1533 static bool equal (omp_declare_variant_base_entry *,
1534 omp_declare_variant_base_entry *);
1535 };
1536
1537 hashval_t
1538 omp_declare_variant_alt_hasher::hash (omp_declare_variant_base_entry *x)
1539 {
1540 return DECL_UID (x->node->decl);
1541 }
1542
1543 bool
1544 omp_declare_variant_alt_hasher::equal (omp_declare_variant_base_entry *x,
1545 omp_declare_variant_base_entry *y)
1546 {
1547 return x->node == y->node;
1548 }
1549
1550 static GTY(()) hash_table<omp_declare_variant_alt_hasher>
1551 *omp_declare_variant_alt;
1552
1553 /* Try to resolve declare variant after gimplification. */
1554
1555 static tree
1556 omp_resolve_late_declare_variant (tree alt)
1557 {
1558 cgraph_node *node = cgraph_node::get (alt);
1559 cgraph_node *cur_node = cgraph_node::get (cfun->decl);
1560 if (node == NULL
1561 || !node->declare_variant_alt
1562 || !cfun->after_inlining)
1563 return alt;
1564
1565 omp_declare_variant_base_entry entry;
1566 entry.base = NULL;
1567 entry.node = node;
1568 entry.variants = NULL;
1569 omp_declare_variant_base_entry *entryp
1570 = omp_declare_variant_alt->find_with_hash (&entry, DECL_UID (alt));
1571
1572 unsigned int i, j;
1573 omp_declare_variant_entry *varentry1, *varentry2;
1574 auto_vec <bool, 16> matches;
1575 unsigned int nmatches = 0;
1576 FOR_EACH_VEC_SAFE_ELT (entryp->variants, i, varentry1)
1577 {
1578 if (varentry1->matches)
1579 {
1580 /* This has been checked to be ok already. */
1581 matches.safe_push (true);
1582 nmatches++;
1583 continue;
1584 }
1585 switch (omp_context_selector_matches (varentry1->ctx))
1586 {
1587 case 0:
1588 matches.safe_push (false);
1589 break;
1590 case -1:
1591 return alt;
1592 default:
1593 matches.safe_push (true);
1594 nmatches++;
1595 break;
1596 }
1597 }
1598
1599 if (nmatches == 0)
1600 return entryp->base->decl;
1601
1602 /* A context selector that is a strict subset of another context selector
1603 has a score of zero. */
1604 FOR_EACH_VEC_SAFE_ELT (entryp->variants, i, varentry1)
1605 if (matches[i])
1606 {
1607 for (j = i + 1;
1608 vec_safe_iterate (entryp->variants, j, &varentry2); ++j)
1609 if (matches[j])
1610 {
1611 int r = omp_context_selector_compare (varentry1->ctx,
1612 varentry2->ctx);
1613 if (r == -1)
1614 {
1615 /* ctx1 is a strict subset of ctx2, ignore ctx1. */
1616 matches[i] = false;
1617 break;
1618 }
1619 else if (r == 1)
1620 /* ctx2 is a strict subset of ctx1, remove ctx2. */
1621 matches[j] = false;
1622 }
1623 }
1624
1625 widest_int max_score = -1;
1626 varentry2 = NULL;
1627 FOR_EACH_VEC_SAFE_ELT (entryp->variants, i, varentry1)
1628 if (matches[i])
1629 {
1630 widest_int score
1631 = (cur_node->simdclone ? varentry1->score_in_declare_simd_clone
1632 : varentry1->score);
1633 if (score > max_score)
1634 {
1635 max_score = score;
1636 varentry2 = varentry1;
1637 }
1638 }
1639 return varentry2->variant->decl;
1640 }
1641
1642 /* Try to resolve declare variant, return the variant decl if it should
1643 be used instead of base, or base otherwise. */
1644
1645 tree
1646 omp_resolve_declare_variant (tree base)
1647 {
1648 tree variant1 = NULL_TREE, variant2 = NULL_TREE;
1649 if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
1650 return omp_resolve_late_declare_variant (base);
1651
1652 auto_vec <tree, 16> variants;
1653 auto_vec <bool, 16> defer;
1654 bool any_deferred = false;
1655 for (tree attr = DECL_ATTRIBUTES (base); attr; attr = TREE_CHAIN (attr))
1656 {
1657 attr = lookup_attribute ("omp declare variant base", attr);
1658 if (attr == NULL_TREE)
1659 break;
1660 if (TREE_CODE (TREE_PURPOSE (TREE_VALUE (attr))) != FUNCTION_DECL)
1661 continue;
1662 switch (omp_context_selector_matches (TREE_VALUE (TREE_VALUE (attr))))
1663 {
1664 case 0:
1665 /* No match, ignore. */
1666 break;
1667 case -1:
1668 /* Needs to be deferred. */
1669 any_deferred = true;
1670 variants.safe_push (attr);
1671 defer.safe_push (true);
1672 break;
1673 default:
1674 variants.safe_push (attr);
1675 defer.safe_push (false);
1676 break;
1677 }
1678 }
1679 if (variants.length () == 0)
1680 return base;
1681
1682 if (any_deferred)
1683 {
1684 widest_int max_score1 = 0;
1685 widest_int max_score2 = 0;
1686 bool first = true;
1687 unsigned int i;
1688 tree attr1, attr2;
1689 omp_declare_variant_base_entry entry;
1690 entry.base = cgraph_node::get_create (base);
1691 entry.node = NULL;
1692 vec_alloc (entry.variants, variants.length ());
1693 FOR_EACH_VEC_ELT (variants, i, attr1)
1694 {
1695 widest_int score1;
1696 widest_int score2;
1697 bool need_two;
1698 tree ctx = TREE_VALUE (TREE_VALUE (attr1));
1699 need_two = omp_context_compute_score (ctx, &score1, false);
1700 if (need_two)
1701 omp_context_compute_score (ctx, &score2, true);
1702 else
1703 score2 = score1;
1704 if (first)
1705 {
1706 first = false;
1707 max_score1 = score1;
1708 max_score2 = score2;
1709 if (!defer[i])
1710 {
1711 variant1 = attr1;
1712 variant2 = attr1;
1713 }
1714 }
1715 else
1716 {
1717 if (max_score1 == score1)
1718 variant1 = NULL_TREE;
1719 else if (score1 > max_score1)
1720 {
1721 max_score1 = score1;
1722 variant1 = defer[i] ? NULL_TREE : attr1;
1723 }
1724 if (max_score2 == score2)
1725 variant2 = NULL_TREE;
1726 else if (score2 > max_score2)
1727 {
1728 max_score2 = score2;
1729 variant2 = defer[i] ? NULL_TREE : attr1;
1730 }
1731 }
1732 omp_declare_variant_entry varentry;
1733 varentry.variant
1734 = cgraph_node::get_create (TREE_PURPOSE (TREE_VALUE (attr1)));
1735 varentry.score = score1;
1736 varentry.score_in_declare_simd_clone = score2;
1737 varentry.ctx = ctx;
1738 varentry.matches = !defer[i];
1739 entry.variants->quick_push (varentry);
1740 }
1741
1742 /* If there is a clear winner variant with the score which is not
1743 deferred, verify it is not a strict subset of any other context
1744 selector and if it is not, it is the best alternative no matter
1745 whether the others do or don't match. */
1746 if (variant1 && variant1 == variant2)
1747 {
1748 tree ctx1 = TREE_VALUE (TREE_VALUE (variant1));
1749 FOR_EACH_VEC_ELT (variants, i, attr2)
1750 {
1751 if (attr2 == variant1)
1752 continue;
1753 tree ctx2 = TREE_VALUE (TREE_VALUE (attr2));
1754 int r = omp_context_selector_compare (ctx1, ctx2);
1755 if (r == -1)
1756 {
1757 /* The winner is a strict subset of ctx2, can't
1758 decide now. */
1759 variant1 = NULL_TREE;
1760 break;
1761 }
1762 }
1763 if (variant1)
1764 {
1765 vec_free (entry.variants);
1766 return TREE_PURPOSE (TREE_VALUE (variant1));
1767 }
1768 }
1769
1770 if (omp_declare_variants == NULL)
1771 omp_declare_variants
1772 = hash_table<omp_declare_variant_hasher>::create_ggc (64);
1773 omp_declare_variant_base_entry **slot
1774 = omp_declare_variants->find_slot (&entry, INSERT);
1775 if (*slot != NULL)
1776 {
1777 vec_free (entry.variants);
1778 return (*slot)->node->decl;
1779 }
1780
1781 *slot = ggc_cleared_alloc<omp_declare_variant_base_entry> ();
1782 (*slot)->base = entry.base;
1783 (*slot)->node = entry.base;
1784 (*slot)->variants = entry.variants;
1785 tree alt = build_decl (DECL_SOURCE_LOCATION (base), FUNCTION_DECL,
1786 DECL_NAME (base), TREE_TYPE (base));
1787 DECL_ARTIFICIAL (alt) = 1;
1788 DECL_IGNORED_P (alt) = 1;
1789 TREE_STATIC (alt) = 1;
1790 tree attributes = DECL_ATTRIBUTES (base);
1791 if (lookup_attribute ("noipa", attributes) == NULL)
1792 {
1793 attributes = tree_cons (get_identifier ("noipa"), NULL, attributes);
1794 if (lookup_attribute ("noinline", attributes) == NULL)
1795 attributes = tree_cons (get_identifier ("noinline"), NULL,
1796 attributes);
1797 if (lookup_attribute ("noclone", attributes) == NULL)
1798 attributes = tree_cons (get_identifier ("noclone"), NULL,
1799 attributes);
1800 if (lookup_attribute ("no_icf", attributes) == NULL)
1801 attributes = tree_cons (get_identifier ("no_icf"), NULL,
1802 attributes);
1803 }
1804 DECL_ATTRIBUTES (alt) = attributes;
1805 DECL_INITIAL (alt) = error_mark_node;
1806 (*slot)->node = cgraph_node::create (alt);
1807 (*slot)->node->declare_variant_alt = 1;
1808 (*slot)->node->create_reference (entry.base, IPA_REF_ADDR);
1809 omp_declare_variant_entry *varentry;
1810 FOR_EACH_VEC_SAFE_ELT (entry.variants, i, varentry)
1811 (*slot)->node->create_reference (varentry->variant, IPA_REF_ADDR);
1812 if (omp_declare_variant_alt == NULL)
1813 omp_declare_variant_alt
1814 = hash_table<omp_declare_variant_alt_hasher>::create_ggc (64);
1815 *omp_declare_variant_alt->find_slot_with_hash (*slot, DECL_UID (alt),
1816 INSERT) = *slot;
1817 return alt;
1818 }
1819
1820 if (variants.length () == 1)
1821 return TREE_PURPOSE (TREE_VALUE (variants[0]));
1822
1823 /* A context selector that is a strict subset of another context selector
1824 has a score of zero. */
1825 tree attr1, attr2;
1826 unsigned int i, j;
1827 FOR_EACH_VEC_ELT (variants, i, attr1)
1828 if (attr1)
1829 {
1830 tree ctx1 = TREE_VALUE (TREE_VALUE (attr1));
1831 FOR_EACH_VEC_ELT_FROM (variants, j, attr2, i + 1)
1832 if (attr2)
1833 {
1834 tree ctx2 = TREE_VALUE (TREE_VALUE (attr2));
1835 int r = omp_context_selector_compare (ctx1, ctx2);
1836 if (r == -1)
1837 {
1838 /* ctx1 is a strict subset of ctx2, remove
1839 attr1 from the vector. */
1840 variants[i] = NULL_TREE;
1841 break;
1842 }
1843 else if (r == 1)
1844 /* ctx2 is a strict subset of ctx1, remove attr2
1845 from the vector. */
1846 variants[j] = NULL_TREE;
1847 }
1848 }
1849 widest_int max_score1 = 0;
1850 widest_int max_score2 = 0;
1851 bool first = true;
1852 FOR_EACH_VEC_ELT (variants, i, attr1)
1853 if (attr1)
1854 {
1855 if (variant1)
1856 {
1857 widest_int score1;
1858 widest_int score2;
1859 bool need_two;
1860 tree ctx;
1861 if (first)
1862 {
1863 first = false;
1864 ctx = TREE_VALUE (TREE_VALUE (variant1));
1865 need_two = omp_context_compute_score (ctx, &max_score1, false);
1866 if (need_two)
1867 omp_context_compute_score (ctx, &max_score2, true);
1868 else
1869 max_score2 = max_score1;
1870 }
1871 ctx = TREE_VALUE (TREE_VALUE (attr1));
1872 need_two = omp_context_compute_score (ctx, &score1, false);
1873 if (need_two)
1874 omp_context_compute_score (ctx, &score2, true);
1875 else
1876 score2 = score1;
1877 if (score1 > max_score1)
1878 {
1879 max_score1 = score1;
1880 variant1 = attr1;
1881 }
1882 if (score2 > max_score2)
1883 {
1884 max_score2 = score2;
1885 variant2 = attr1;
1886 }
1887 }
1888 else
1889 {
1890 variant1 = attr1;
1891 variant2 = attr1;
1892 }
1893 }
1894 /* If there is a disagreement on which variant has the highest score
1895 depending on whether it will be in a declare simd clone or not,
1896 punt for now and defer until after IPA where we will know that. */
1897 return ((variant1 && variant1 == variant2)
1898 ? TREE_PURPOSE (TREE_VALUE (variant1)) : base);
1899 }
1900
1901
1902 /* Encode an oacc launch argument. This matches the GOMP_LAUNCH_PACK
1903 macro on gomp-constants.h. We do not check for overflow. */
1904
1905 tree
1906 oacc_launch_pack (unsigned code, tree device, unsigned op)
1907 {
1908 tree res;
1909
1910 res = build_int_cst (unsigned_type_node, GOMP_LAUNCH_PACK (code, 0, op));
1911 if (device)
1912 {
1913 device = fold_build2 (LSHIFT_EXPR, unsigned_type_node,
1914 device, build_int_cst (unsigned_type_node,
1915 GOMP_LAUNCH_DEVICE_SHIFT));
1916 res = fold_build2 (BIT_IOR_EXPR, unsigned_type_node, res, device);
1917 }
1918 return res;
1919 }
1920
1921 /* FIXME: What is the following comment for? */
1922 /* Look for compute grid dimension clauses and convert to an attribute
1923 attached to FN. This permits the target-side code to (a) massage
1924 the dimensions, (b) emit that data and (c) optimize. Non-constant
1925 dimensions are pushed onto ARGS.
1926
1927 The attribute value is a TREE_LIST. A set of dimensions is
1928 represented as a list of INTEGER_CST. Those that are runtime
1929 exprs are represented as an INTEGER_CST of zero.
1930
1931 TODO: Normally the attribute will just contain a single such list. If
1932 however it contains a list of lists, this will represent the use of
1933 device_type. Each member of the outer list is an assoc list of
1934 dimensions, keyed by the device type. The first entry will be the
1935 default. Well, that's the plan. */
1936
1937 /* Replace any existing oacc fn attribute with updated dimensions. */
1938
1939 /* Variant working on a list of attributes. */
1940
1941 tree
1942 oacc_replace_fn_attrib_attr (tree attribs, tree dims)
1943 {
1944 tree ident = get_identifier (OACC_FN_ATTRIB);
1945
1946 /* If we happen to be present as the first attrib, drop it. */
1947 if (attribs && TREE_PURPOSE (attribs) == ident)
1948 attribs = TREE_CHAIN (attribs);
1949 return tree_cons (ident, dims, attribs);
1950 }
1951
1952 /* Variant working on a function decl. */
1953
1954 void
1955 oacc_replace_fn_attrib (tree fn, tree dims)
1956 {
1957 DECL_ATTRIBUTES (fn)
1958 = oacc_replace_fn_attrib_attr (DECL_ATTRIBUTES (fn), dims);
1959 }
1960
1961 /* Scan CLAUSES for launch dimensions and attach them to the oacc
1962 function attribute. Push any that are non-constant onto the ARGS
1963 list, along with an appropriate GOMP_LAUNCH_DIM tag. */
1964
1965 void
1966 oacc_set_fn_attrib (tree fn, tree clauses, vec<tree> *args)
1967 {
1968 /* Must match GOMP_DIM ordering. */
1969 static const omp_clause_code ids[]
1970 = { OMP_CLAUSE_NUM_GANGS, OMP_CLAUSE_NUM_WORKERS,
1971 OMP_CLAUSE_VECTOR_LENGTH };
1972 unsigned ix;
1973 tree dims[GOMP_DIM_MAX];
1974
1975 tree attr = NULL_TREE;
1976 unsigned non_const = 0;
1977
1978 for (ix = GOMP_DIM_MAX; ix--;)
1979 {
1980 tree clause = omp_find_clause (clauses, ids[ix]);
1981 tree dim = NULL_TREE;
1982
1983 if (clause)
1984 dim = OMP_CLAUSE_EXPR (clause, ids[ix]);
1985 dims[ix] = dim;
1986 if (dim && TREE_CODE (dim) != INTEGER_CST)
1987 {
1988 dim = integer_zero_node;
1989 non_const |= GOMP_DIM_MASK (ix);
1990 }
1991 attr = tree_cons (NULL_TREE, dim, attr);
1992 }
1993
1994 oacc_replace_fn_attrib (fn, attr);
1995
1996 if (non_const)
1997 {
1998 /* Push a dynamic argument set. */
1999 args->safe_push (oacc_launch_pack (GOMP_LAUNCH_DIM,
2000 NULL_TREE, non_const));
2001 for (unsigned ix = 0; ix != GOMP_DIM_MAX; ix++)
2002 if (non_const & GOMP_DIM_MASK (ix))
2003 args->safe_push (dims[ix]);
2004 }
2005 }
2006
2007 /* Verify OpenACC routine clauses.
2008
2009 Returns 0 if FNDECL should be marked with an OpenACC 'routine' directive, 1
2010 if it has already been marked in compatible way, and -1 if incompatible.
2011 Upon returning, the chain of clauses will contain exactly one clause
2012 specifying the level of parallelism. */
2013
2014 int
2015 oacc_verify_routine_clauses (tree fndecl, tree *clauses, location_t loc,
2016 const char *routine_str)
2017 {
2018 tree c_level = NULL_TREE;
2019 tree c_p = NULL_TREE;
2020 for (tree c = *clauses; c; c_p = c, c = OMP_CLAUSE_CHAIN (c))
2021 switch (OMP_CLAUSE_CODE (c))
2022 {
2023 case OMP_CLAUSE_GANG:
2024 case OMP_CLAUSE_WORKER:
2025 case OMP_CLAUSE_VECTOR:
2026 case OMP_CLAUSE_SEQ:
2027 if (c_level == NULL_TREE)
2028 c_level = c;
2029 else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_CODE (c_level))
2030 {
2031 /* This has already been diagnosed in the front ends. */
2032 /* Drop the duplicate clause. */
2033 gcc_checking_assert (c_p != NULL_TREE);
2034 OMP_CLAUSE_CHAIN (c_p) = OMP_CLAUSE_CHAIN (c);
2035 c = c_p;
2036 }
2037 else
2038 {
2039 error_at (OMP_CLAUSE_LOCATION (c),
2040 "%qs specifies a conflicting level of parallelism",
2041 omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
2042 inform (OMP_CLAUSE_LOCATION (c_level),
2043 "... to the previous %qs clause here",
2044 omp_clause_code_name[OMP_CLAUSE_CODE (c_level)]);
2045 /* Drop the conflicting clause. */
2046 gcc_checking_assert (c_p != NULL_TREE);
2047 OMP_CLAUSE_CHAIN (c_p) = OMP_CLAUSE_CHAIN (c);
2048 c = c_p;
2049 }
2050 break;
2051 default:
2052 gcc_unreachable ();
2053 }
2054 if (c_level == NULL_TREE)
2055 {
2056 /* Default to an implicit 'seq' clause. */
2057 c_level = build_omp_clause (loc, OMP_CLAUSE_SEQ);
2058 OMP_CLAUSE_CHAIN (c_level) = *clauses;
2059 *clauses = c_level;
2060 }
2061 /* In *clauses, we now have exactly one clause specifying the level of
2062 parallelism. */
2063
2064 tree attr
2065 = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (fndecl));
2066 if (attr != NULL_TREE)
2067 {
2068 /* Diagnose if "#pragma omp declare target" has also been applied. */
2069 if (TREE_VALUE (attr) == NULL_TREE)
2070 {
2071 /* See <https://gcc.gnu.org/PR93465>; the semantics of combining
2072 OpenACC and OpenMP 'target' are not clear. */
2073 error_at (loc,
2074 "cannot apply %<%s%> to %qD, which has also been"
2075 " marked with an OpenMP 'declare target' directive",
2076 routine_str, fndecl);
2077 /* Incompatible. */
2078 return -1;
2079 }
2080
2081 /* If a "#pragma acc routine" has already been applied, just verify
2082 this one for compatibility. */
2083 /* Collect previous directive's clauses. */
2084 tree c_level_p = NULL_TREE;
2085 for (tree c = TREE_VALUE (attr); c; c = OMP_CLAUSE_CHAIN (c))
2086 switch (OMP_CLAUSE_CODE (c))
2087 {
2088 case OMP_CLAUSE_GANG:
2089 case OMP_CLAUSE_WORKER:
2090 case OMP_CLAUSE_VECTOR:
2091 case OMP_CLAUSE_SEQ:
2092 gcc_checking_assert (c_level_p == NULL_TREE);
2093 c_level_p = c;
2094 break;
2095 default:
2096 gcc_unreachable ();
2097 }
2098 gcc_checking_assert (c_level_p != NULL_TREE);
2099 /* ..., and compare to current directive's, which we've already collected
2100 above. */
2101 tree c_diag;
2102 tree c_diag_p;
2103 /* Matching level of parallelism? */
2104 if (OMP_CLAUSE_CODE (c_level) != OMP_CLAUSE_CODE (c_level_p))
2105 {
2106 c_diag = c_level;
2107 c_diag_p = c_level_p;
2108 goto incompatible;
2109 }
2110 /* Compatible. */
2111 return 1;
2112
2113 incompatible:
2114 if (c_diag != NULL_TREE)
2115 error_at (OMP_CLAUSE_LOCATION (c_diag),
2116 "incompatible %qs clause when applying"
2117 " %<%s%> to %qD, which has already been"
2118 " marked with an OpenACC 'routine' directive",
2119 omp_clause_code_name[OMP_CLAUSE_CODE (c_diag)],
2120 routine_str, fndecl);
2121 else if (c_diag_p != NULL_TREE)
2122 error_at (loc,
2123 "missing %qs clause when applying"
2124 " %<%s%> to %qD, which has already been"
2125 " marked with an OpenACC 'routine' directive",
2126 omp_clause_code_name[OMP_CLAUSE_CODE (c_diag_p)],
2127 routine_str, fndecl);
2128 else
2129 gcc_unreachable ();
2130 if (c_diag_p != NULL_TREE)
2131 inform (OMP_CLAUSE_LOCATION (c_diag_p),
2132 "... with %qs clause here",
2133 omp_clause_code_name[OMP_CLAUSE_CODE (c_diag_p)]);
2134 else
2135 {
2136 /* In the front ends, we don't preserve location information for the
2137 OpenACC routine directive itself. However, that of c_level_p
2138 should be close. */
2139 location_t loc_routine = OMP_CLAUSE_LOCATION (c_level_p);
2140 inform (loc_routine, "... without %qs clause near to here",
2141 omp_clause_code_name[OMP_CLAUSE_CODE (c_diag)]);
2142 }
2143 /* Incompatible. */
2144 return -1;
2145 }
2146
2147 return 0;
2148 }
2149
2150 /* Process the OpenACC 'routine' directive clauses to generate an attribute
2151 for the level of parallelism. All dimensions have a size of zero
2152 (dynamic). TREE_PURPOSE is set to indicate whether that dimension
2153 can have a loop partitioned on it. non-zero indicates
2154 yes, zero indicates no. By construction once a non-zero has been
2155 reached, further inner dimensions must also be non-zero. We set
2156 TREE_VALUE to zero for the dimensions that may be partitioned and
2157 1 for the other ones -- if a loop is (erroneously) spawned at
2158 an outer level, we don't want to try and partition it. */
2159
2160 tree
2161 oacc_build_routine_dims (tree clauses)
2162 {
2163 /* Must match GOMP_DIM ordering. */
2164 static const omp_clause_code ids[]
2165 = {OMP_CLAUSE_GANG, OMP_CLAUSE_WORKER, OMP_CLAUSE_VECTOR, OMP_CLAUSE_SEQ};
2166 int ix;
2167 int level = -1;
2168
2169 for (; clauses; clauses = OMP_CLAUSE_CHAIN (clauses))
2170 for (ix = GOMP_DIM_MAX + 1; ix--;)
2171 if (OMP_CLAUSE_CODE (clauses) == ids[ix])
2172 {
2173 level = ix;
2174 break;
2175 }
2176 gcc_checking_assert (level >= 0);
2177
2178 tree dims = NULL_TREE;
2179
2180 for (ix = GOMP_DIM_MAX; ix--;)
2181 dims = tree_cons (build_int_cst (boolean_type_node, ix >= level),
2182 build_int_cst (integer_type_node, ix < level), dims);
2183
2184 return dims;
2185 }
2186
2187 /* Retrieve the oacc function attrib and return it. Non-oacc
2188 functions will return NULL. */
2189
2190 tree
2191 oacc_get_fn_attrib (tree fn)
2192 {
2193 return lookup_attribute (OACC_FN_ATTRIB, DECL_ATTRIBUTES (fn));
2194 }
2195
2196 /* Return true if FN is an OpenMP or OpenACC offloading function. */
2197
2198 bool
2199 offloading_function_p (tree fn)
2200 {
2201 tree attrs = DECL_ATTRIBUTES (fn);
2202 return (lookup_attribute ("omp declare target", attrs)
2203 || lookup_attribute ("omp target entrypoint", attrs));
2204 }
2205
2206 /* Extract an oacc execution dimension from FN. FN must be an
2207 offloaded function or routine that has already had its execution
2208 dimensions lowered to the target-specific values. */
2209
2210 int
2211 oacc_get_fn_dim_size (tree fn, int axis)
2212 {
2213 tree attrs = oacc_get_fn_attrib (fn);
2214
2215 gcc_assert (axis < GOMP_DIM_MAX);
2216
2217 tree dims = TREE_VALUE (attrs);
2218 while (axis--)
2219 dims = TREE_CHAIN (dims);
2220
2221 int size = TREE_INT_CST_LOW (TREE_VALUE (dims));
2222
2223 return size;
2224 }
2225
2226 /* Extract the dimension axis from an IFN_GOACC_DIM_POS or
2227 IFN_GOACC_DIM_SIZE call. */
2228
2229 int
2230 oacc_get_ifn_dim_arg (const gimple *stmt)
2231 {
2232 gcc_checking_assert (gimple_call_internal_fn (stmt) == IFN_GOACC_DIM_SIZE
2233 || gimple_call_internal_fn (stmt) == IFN_GOACC_DIM_POS);
2234 tree arg = gimple_call_arg (stmt, 0);
2235 HOST_WIDE_INT axis = TREE_INT_CST_LOW (arg);
2236
2237 gcc_checking_assert (axis >= 0 && axis < GOMP_DIM_MAX);
2238 return (int) axis;
2239 }
2240
2241 #include "gt-omp-general.h"