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