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 #ifdef ACCEL_COMPILER
936 return 0;
937 #else
938 if (omp_maybe_offloaded ())
939 ret = -1;
940 continue;
941 #endif
942 }
943 if (!strcmp (prop, "nohost"))
944 {
945 #ifndef ACCEL_COMPILER
946 if (omp_maybe_offloaded ())
947 ret = -1;
948 else
949 return 0;
950 #endif
951 continue;
952 }
953 int r = 0;
954 if (targetm.omp.device_kind_arch_isa != NULL)
955 r = targetm.omp.device_kind_arch_isa (omp_device_kind,
956 prop);
957 else
958 r = strcmp (prop, "cpu") == 0;
959 if (r == 0 || (r == -1 && symtab->state != PARSING))
960 {
961 /* If we are or might be in a target region or
962 declare target function, need to take into account
963 also offloading values. */
964 if (!omp_maybe_offloaded ())
965 return 0;
966 if (strcmp (prop, "gpu") == 0
967 && hsa_gen_requested_p ())
968 {
969 ret = -1;
970 continue;
971 }
972 if (ENABLE_OFFLOADING)
973 {
974 const char *kinds = omp_offload_device_kind;
975 if (omp_offload_device_kind_arch_isa (kinds, prop))
976 {
977 ret = -1;
978 continue;
979 }
980 }
981 return 0;
982 }
983 else if (r == -1)
984 ret = -1;
985 /* If kind matches on the host, it still might not match
986 in the offloading region. */
987 else if (omp_maybe_offloaded ())
988 ret = -1;
989 }
990 break;
991 case 'i':
992 if (set == 'd' && !strcmp (sel, "isa"))
993 for (tree t3 = TREE_VALUE (t2); t3; t3 = TREE_CHAIN (t3))
994 {
995 const char *isa = omp_context_name_list_prop (t3);
996 if (isa == NULL)
997 return 0;
998 int r = 0;
999 if (targetm.omp.device_kind_arch_isa != NULL)
1000 r = targetm.omp.device_kind_arch_isa (omp_device_isa,
1001 isa);
1002 if (r == 0 || (r == -1 && symtab->state != PARSING))
1003 {
1004 /* If isa is valid on the target, but not in the
1005 current function and current function has
1006 #pragma omp declare simd on it, some simd clones
1007 might have the isa added later on. */
1008 if (r == -1
1009 && targetm.simd_clone.compute_vecsize_and_simdlen)
1010 {
1011 tree attrs
1012 = DECL_ATTRIBUTES (current_function_decl);
1013 if (lookup_attribute ("omp declare simd", attrs))
1014 {
1015 ret = -1;
1016 continue;
1017 }
1018 }
1019 /* If we are or might be in a target region or
1020 declare target function, need to take into account
1021 also offloading values. */
1022 if (!omp_maybe_offloaded ())
1023 return 0;
1024 if (ENABLE_OFFLOADING)
1025 {
1026 const char *isas = omp_offload_device_isa;
1027 if (omp_offload_device_kind_arch_isa (isas, isa))
1028 {
1029 ret = -1;
1030 continue;
1031 }
1032 }
1033 return 0;
1034 }
1035 else if (r == -1)
1036 ret = -1;
1037 /* If isa matches on the host, it still might not match
1038 in the offloading region. */
1039 else if (omp_maybe_offloaded ())
1040 ret = -1;
1041 }
1042 break;
1043 case 'c':
1044 if (set == 'u' && !strcmp (sel, "condition"))
1045 for (tree t3 = TREE_VALUE (t2); t3; t3 = TREE_CHAIN (t3))
1046 if (TREE_PURPOSE (t3) == NULL_TREE)
1047 {
1048 if (integer_zerop (TREE_VALUE (t3)))
1049 return 0;
1050 if (integer_nonzerop (TREE_VALUE (t3)))
1051 break;
1052 ret = -1;
1053 }
1054 break;
1055 default:
1056 break;
1057 }
1058 }
1059 }
1060 return ret;
1061 }
1062
1063 /* Compare construct={simd} CLAUSES1 with CLAUSES2, return 0/-1/1/2 as
1064 in omp_context_selector_set_compare. */
1065
1066 static int
omp_construct_simd_compare(tree clauses1,tree clauses2)1067 omp_construct_simd_compare (tree clauses1, tree clauses2)
1068 {
1069 if (clauses1 == NULL_TREE)
1070 return clauses2 == NULL_TREE ? 0 : -1;
1071 if (clauses2 == NULL_TREE)
1072 return 1;
1073
1074 int r = 0;
1075 struct declare_variant_simd_data {
1076 bool inbranch, notinbranch;
1077 tree simdlen;
1078 auto_vec<tree,16> data_sharing;
1079 auto_vec<tree,16> aligned;
1080 declare_variant_simd_data ()
1081 : inbranch(false), notinbranch(false), simdlen(NULL_TREE) {}
1082 } data[2];
1083 unsigned int i;
1084 for (i = 0; i < 2; i++)
1085 for (tree c = i ? clauses2 : clauses1; c; c = OMP_CLAUSE_CHAIN (c))
1086 {
1087 vec<tree> *v;
1088 switch (OMP_CLAUSE_CODE (c))
1089 {
1090 case OMP_CLAUSE_INBRANCH:
1091 data[i].inbranch = true;
1092 continue;
1093 case OMP_CLAUSE_NOTINBRANCH:
1094 data[i].notinbranch = true;
1095 continue;
1096 case OMP_CLAUSE_SIMDLEN:
1097 data[i].simdlen = OMP_CLAUSE_SIMDLEN_EXPR (c);
1098 continue;
1099 case OMP_CLAUSE_UNIFORM:
1100 case OMP_CLAUSE_LINEAR:
1101 v = &data[i].data_sharing;
1102 break;
1103 case OMP_CLAUSE_ALIGNED:
1104 v = &data[i].aligned;
1105 break;
1106 default:
1107 gcc_unreachable ();
1108 }
1109 unsigned HOST_WIDE_INT argno = tree_to_uhwi (OMP_CLAUSE_DECL (c));
1110 if (argno >= v->length ())
1111 v->safe_grow_cleared (argno + 1);
1112 (*v)[argno] = c;
1113 }
1114 /* Here, r is used as a bitmask, 2 is set if CLAUSES1 has something
1115 CLAUSES2 doesn't, 1 is set if CLAUSES2 has something CLAUSES1
1116 doesn't. Thus, r == 3 implies return value 2, r == 1 implies
1117 -1, r == 2 implies 1 and r == 0 implies 0. */
1118 if (data[0].inbranch != data[1].inbranch)
1119 r |= data[0].inbranch ? 2 : 1;
1120 if (data[0].notinbranch != data[1].notinbranch)
1121 r |= data[0].notinbranch ? 2 : 1;
1122 if (!simple_cst_equal (data[0].simdlen, data[1].simdlen))
1123 {
1124 if (data[0].simdlen && data[1].simdlen)
1125 return 2;
1126 r |= data[0].simdlen ? 2 : 1;
1127 }
1128 if (data[0].data_sharing.length () < data[1].data_sharing.length ()
1129 || data[0].aligned.length () < data[1].aligned.length ())
1130 r |= 1;
1131 tree c1, c2;
1132 FOR_EACH_VEC_ELT (data[0].data_sharing, i, c1)
1133 {
1134 c2 = (i < data[1].data_sharing.length ()
1135 ? data[1].data_sharing[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 (OMP_CLAUSE_CODE (c1) != OMP_CLAUSE_CODE (c2))
1144 return 2;
1145 if (OMP_CLAUSE_CODE (c1) != OMP_CLAUSE_LINEAR)
1146 continue;
1147 if (OMP_CLAUSE_LINEAR_VARIABLE_STRIDE (c1)
1148 != OMP_CLAUSE_LINEAR_VARIABLE_STRIDE (c2))
1149 return 2;
1150 if (OMP_CLAUSE_LINEAR_KIND (c1) != OMP_CLAUSE_LINEAR_KIND (c2))
1151 return 2;
1152 if (!simple_cst_equal (OMP_CLAUSE_LINEAR_STEP (c1),
1153 OMP_CLAUSE_LINEAR_STEP (c2)))
1154 return 2;
1155 }
1156 FOR_EACH_VEC_ELT (data[0].aligned, i, c1)
1157 {
1158 c2 = i < data[1].aligned.length () ? data[1].aligned[i] : NULL_TREE;
1159 if ((c1 == NULL_TREE) != (c2 == NULL_TREE))
1160 {
1161 r |= c1 != NULL_TREE ? 2 : 1;
1162 continue;
1163 }
1164 if (c1 == NULL_TREE)
1165 continue;
1166 if (!simple_cst_equal (OMP_CLAUSE_ALIGNED_ALIGNMENT (c1),
1167 OMP_CLAUSE_ALIGNED_ALIGNMENT (c2)))
1168 return 2;
1169 }
1170 switch (r)
1171 {
1172 case 0: return 0;
1173 case 1: return -1;
1174 case 2: return 1;
1175 case 3: return 2;
1176 default: gcc_unreachable ();
1177 }
1178 }
1179
1180 /* Compare properties of selectors SEL from SET other than construct.
1181 Return 0/-1/1/2 as in omp_context_selector_set_compare.
1182 Unlike set names or selector names, properties can have duplicates. */
1183
1184 static int
omp_context_selector_props_compare(const char * set,const char * sel,tree ctx1,tree ctx2)1185 omp_context_selector_props_compare (const char *set, const char *sel,
1186 tree ctx1, tree ctx2)
1187 {
1188 int ret = 0;
1189 for (int pass = 0; pass < 2; pass++)
1190 for (tree t1 = pass ? ctx2 : ctx1; t1; t1 = TREE_CHAIN (t1))
1191 {
1192 tree t2;
1193 for (t2 = pass ? ctx1 : ctx2; t2; t2 = TREE_CHAIN (t2))
1194 if (TREE_PURPOSE (t1) == TREE_PURPOSE (t2))
1195 {
1196 if (TREE_PURPOSE (t1) == NULL_TREE)
1197 {
1198 if (set[0] == 'u' && strcmp (sel, "condition") == 0)
1199 {
1200 if (integer_zerop (TREE_VALUE (t1))
1201 != integer_zerop (TREE_VALUE (t2)))
1202 return 2;
1203 break;
1204 }
1205 if (simple_cst_equal (TREE_VALUE (t1), TREE_VALUE (t2)))
1206 break;
1207 }
1208 else if (strcmp (IDENTIFIER_POINTER (TREE_PURPOSE (t1)),
1209 " score") == 0)
1210 {
1211 if (!simple_cst_equal (TREE_VALUE (t1), TREE_VALUE (t2)))
1212 return 2;
1213 break;
1214 }
1215 else
1216 break;
1217 }
1218 else if (TREE_PURPOSE (t1)
1219 && TREE_PURPOSE (t2) == NULL_TREE
1220 && TREE_CODE (TREE_VALUE (t2)) == STRING_CST)
1221 {
1222 const char *p1 = omp_context_name_list_prop (t1);
1223 const char *p2 = omp_context_name_list_prop (t2);
1224 if (p2
1225 && strcmp (p1, p2) == 0
1226 && strcmp (p1, " score"))
1227 break;
1228 }
1229 else if (TREE_PURPOSE (t1) == NULL_TREE
1230 && TREE_PURPOSE (t2)
1231 && TREE_CODE (TREE_VALUE (t1)) == STRING_CST)
1232 {
1233 const char *p1 = omp_context_name_list_prop (t1);
1234 const char *p2 = omp_context_name_list_prop (t2);
1235 if (p1
1236 && strcmp (p1, p2) == 0
1237 && strcmp (p1, " score"))
1238 break;
1239 }
1240 if (t2 == NULL_TREE)
1241 {
1242 int r = pass ? -1 : 1;
1243 if (ret && ret != r)
1244 return 2;
1245 else if (pass)
1246 return r;
1247 else
1248 {
1249 ret = r;
1250 break;
1251 }
1252 }
1253 }
1254 return ret;
1255 }
1256
1257 /* Compare single context selector sets CTX1 and CTX2 with SET name.
1258 Return 0 if CTX1 is equal to CTX2,
1259 -1 if CTX1 is a strict subset of CTX2,
1260 1 if CTX2 is a strict subset of CTX1, or
1261 2 if neither context is a subset of another one. */
1262
1263 int
omp_context_selector_set_compare(const char * set,tree ctx1,tree ctx2)1264 omp_context_selector_set_compare (const char *set, tree ctx1, tree ctx2)
1265 {
1266 bool swapped = false;
1267 int ret = 0;
1268 int len1 = list_length (ctx1);
1269 int len2 = list_length (ctx2);
1270 int cnt = 0;
1271 if (len1 < len2)
1272 {
1273 swapped = true;
1274 std::swap (ctx1, ctx2);
1275 std::swap (len1, len2);
1276 }
1277 if (set[0] == 'c')
1278 {
1279 tree t1;
1280 tree t2 = ctx2;
1281 tree simd = get_identifier ("simd");
1282 /* Handle construct set specially. In this case the order
1283 of the selector matters too. */
1284 for (t1 = ctx1; t1; t1 = TREE_CHAIN (t1))
1285 if (TREE_PURPOSE (t1) == TREE_PURPOSE (t2))
1286 {
1287 int r = 0;
1288 if (TREE_PURPOSE (t1) == simd)
1289 r = omp_construct_simd_compare (TREE_VALUE (t1),
1290 TREE_VALUE (t2));
1291 if (r == 2 || (ret && r && (ret < 0) != (r < 0)))
1292 return 2;
1293 if (ret == 0)
1294 ret = r;
1295 t2 = TREE_CHAIN (t2);
1296 if (t2 == NULL_TREE)
1297 {
1298 t1 = TREE_CHAIN (t1);
1299 break;
1300 }
1301 }
1302 else if (ret < 0)
1303 return 2;
1304 else
1305 ret = 1;
1306 if (t2 != NULL_TREE)
1307 return 2;
1308 if (t1 != NULL_TREE)
1309 {
1310 if (ret < 0)
1311 return 2;
1312 ret = 1;
1313 }
1314 if (ret == 0)
1315 return 0;
1316 return swapped ? -ret : ret;
1317 }
1318 for (tree t1 = ctx1; t1; t1 = TREE_CHAIN (t1))
1319 {
1320 tree t2;
1321 for (t2 = ctx2; t2; t2 = TREE_CHAIN (t2))
1322 if (TREE_PURPOSE (t1) == TREE_PURPOSE (t2))
1323 {
1324 const char *sel = IDENTIFIER_POINTER (TREE_PURPOSE (t1));
1325 int r = omp_context_selector_props_compare (set, sel,
1326 TREE_VALUE (t1),
1327 TREE_VALUE (t2));
1328 if (r == 2 || (ret && r && (ret < 0) != (r < 0)))
1329 return 2;
1330 if (ret == 0)
1331 ret = r;
1332 cnt++;
1333 break;
1334 }
1335 if (t2 == NULL_TREE)
1336 {
1337 if (ret == -1)
1338 return 2;
1339 ret = 1;
1340 }
1341 }
1342 if (cnt < len2)
1343 return 2;
1344 if (ret == 0)
1345 return 0;
1346 return swapped ? -ret : ret;
1347 }
1348
1349 /* Compare whole context selector specification CTX1 and CTX2.
1350 Return 0 if CTX1 is equal to CTX2,
1351 -1 if CTX1 is a strict subset of CTX2,
1352 1 if CTX2 is a strict subset of CTX1, or
1353 2 if neither context is a subset of another one. */
1354
1355 static int
omp_context_selector_compare(tree ctx1,tree ctx2)1356 omp_context_selector_compare (tree ctx1, tree ctx2)
1357 {
1358 bool swapped = false;
1359 int ret = 0;
1360 int len1 = list_length (ctx1);
1361 int len2 = list_length (ctx2);
1362 int cnt = 0;
1363 if (len1 < len2)
1364 {
1365 swapped = true;
1366 std::swap (ctx1, ctx2);
1367 std::swap (len1, len2);
1368 }
1369 for (tree t1 = ctx1; t1; t1 = TREE_CHAIN (t1))
1370 {
1371 tree t2;
1372 for (t2 = ctx2; t2; t2 = TREE_CHAIN (t2))
1373 if (TREE_PURPOSE (t1) == TREE_PURPOSE (t2))
1374 {
1375 const char *set = IDENTIFIER_POINTER (TREE_PURPOSE (t1));
1376 int r = omp_context_selector_set_compare (set, TREE_VALUE (t1),
1377 TREE_VALUE (t2));
1378 if (r == 2 || (ret && r && (ret < 0) != (r < 0)))
1379 return 2;
1380 if (ret == 0)
1381 ret = r;
1382 cnt++;
1383 break;
1384 }
1385 if (t2 == NULL_TREE)
1386 {
1387 if (ret == -1)
1388 return 2;
1389 ret = 1;
1390 }
1391 }
1392 if (cnt < len2)
1393 return 2;
1394 if (ret == 0)
1395 return 0;
1396 return swapped ? -ret : ret;
1397 }
1398
1399 /* From context selector CTX, return trait-selector with name SEL in
1400 trait-selector-set with name SET if any, or NULL_TREE if not found.
1401 If SEL is NULL, return the list of trait-selectors in SET. */
1402
1403 tree
omp_get_context_selector(tree ctx,const char * set,const char * sel)1404 omp_get_context_selector (tree ctx, const char *set, const char *sel)
1405 {
1406 tree setid = get_identifier (set);
1407 tree selid = sel ? get_identifier (sel) : NULL_TREE;
1408 for (tree t1 = ctx; t1; t1 = TREE_CHAIN (t1))
1409 if (TREE_PURPOSE (t1) == setid)
1410 {
1411 if (sel == NULL)
1412 return TREE_VALUE (t1);
1413 for (tree t2 = TREE_VALUE (t1); t2; t2 = TREE_CHAIN (t2))
1414 if (TREE_PURPOSE (t2) == selid)
1415 return t2;
1416 }
1417 return NULL_TREE;
1418 }
1419
1420 /* Compute *SCORE for context selector CTX. Return true if the score
1421 would be different depending on whether it is a declare simd clone or
1422 not. DECLARE_SIMD should be true for the case when it would be
1423 a declare simd clone. */
1424
1425 static bool
omp_context_compute_score(tree ctx,widest_int * score,bool declare_simd)1426 omp_context_compute_score (tree ctx, widest_int *score, bool declare_simd)
1427 {
1428 tree construct = omp_get_context_selector (ctx, "construct", NULL);
1429 bool has_kind = omp_get_context_selector (ctx, "device", "kind");
1430 bool has_arch = omp_get_context_selector (ctx, "device", "arch");
1431 bool has_isa = omp_get_context_selector (ctx, "device", "isa");
1432 bool ret = false;
1433 *score = 1;
1434 for (tree t1 = ctx; t1; t1 = TREE_CHAIN (t1))
1435 if (TREE_VALUE (t1) != construct)
1436 for (tree t2 = TREE_VALUE (t1); t2; t2 = TREE_CHAIN (t2))
1437 if (tree t3 = TREE_VALUE (t2))
1438 if (TREE_PURPOSE (t3)
1439 && strcmp (IDENTIFIER_POINTER (TREE_PURPOSE (t3)), " score") == 0
1440 && TREE_CODE (TREE_VALUE (t3)) == INTEGER_CST)
1441 *score += wi::to_widest (TREE_VALUE (t3));
1442 if (construct || has_kind || has_arch || has_isa)
1443 {
1444 int scores[12];
1445 enum tree_code constructs[5];
1446 int nconstructs = 0;
1447 if (construct)
1448 nconstructs = omp_constructor_traits_to_codes (construct, constructs);
1449 if (omp_construct_selector_matches (constructs, nconstructs, scores)
1450 == 2)
1451 ret = true;
1452 int b = declare_simd ? nconstructs + 1 : 0;
1453 if (scores[b + nconstructs] + 4U < score->get_precision ())
1454 {
1455 for (int n = 0; n < nconstructs; ++n)
1456 {
1457 if (scores[b + n] < 0)
1458 {
1459 *score = -1;
1460 return ret;
1461 }
1462 *score += wi::shifted_mask <widest_int> (scores[b + n], 1, false);
1463 }
1464 if (has_kind)
1465 *score += wi::shifted_mask <widest_int> (scores[b + nconstructs],
1466 1, false);
1467 if (has_arch)
1468 *score += wi::shifted_mask <widest_int> (scores[b + nconstructs] + 1,
1469 1, false);
1470 if (has_isa)
1471 *score += wi::shifted_mask <widest_int> (scores[b + nconstructs] + 2,
1472 1, false);
1473 }
1474 else /* FIXME: Implement this. */
1475 gcc_unreachable ();
1476 }
1477 return ret;
1478 }
1479
1480 /* Try to resolve declare variant, return the variant decl if it should
1481 be used instead of base, or base otherwise. */
1482
1483 tree
omp_resolve_declare_variant(tree base)1484 omp_resolve_declare_variant (tree base)
1485 {
1486 tree variant1 = NULL_TREE, variant2 = NULL_TREE;
1487 auto_vec <tree, 16> variants;
1488 auto_vec <bool, 16> defer;
1489 bool any_deferred = false;
1490 for (tree attr = DECL_ATTRIBUTES (base); attr; attr = TREE_CHAIN (attr))
1491 {
1492 attr = lookup_attribute ("omp declare variant base", attr);
1493 if (attr == NULL_TREE)
1494 break;
1495 if (TREE_CODE (TREE_PURPOSE (TREE_VALUE (attr))) != FUNCTION_DECL)
1496 continue;
1497 switch (omp_context_selector_matches (TREE_VALUE (TREE_VALUE (attr))))
1498 {
1499 case 0:
1500 /* No match, ignore. */
1501 break;
1502 case -1:
1503 /* Needs to be deferred. */
1504 any_deferred = true;
1505 variants.safe_push (attr);
1506 defer.safe_push (true);
1507 break;
1508 default:
1509 variants.safe_push (attr);
1510 defer.safe_push (false);
1511 break;
1512 }
1513 }
1514 if (variants.length () == 0)
1515 return base;
1516
1517 if (any_deferred)
1518 {
1519 widest_int max_score1 = 0;
1520 widest_int max_score2 = 0;
1521 bool first = true;
1522 unsigned int i;
1523 tree attr1, attr2;
1524 FOR_EACH_VEC_ELT (variants, i, attr1)
1525 {
1526 widest_int score1;
1527 widest_int score2;
1528 bool need_two;
1529 tree ctx = TREE_VALUE (TREE_VALUE (attr1));
1530 need_two = omp_context_compute_score (ctx, &score1, false);
1531 if (need_two)
1532 omp_context_compute_score (ctx, &score2, true);
1533 else
1534 score2 = score1;
1535 if (first)
1536 {
1537 first = false;
1538 max_score1 = score1;
1539 max_score2 = score2;
1540 if (!defer[i])
1541 {
1542 variant1 = attr1;
1543 variant2 = attr1;
1544 }
1545 }
1546 else
1547 {
1548 if (max_score1 == score1)
1549 variant1 = NULL_TREE;
1550 else if (score1 > max_score1)
1551 {
1552 max_score1 = score1;
1553 variant1 = defer[i] ? NULL_TREE : attr1;
1554 }
1555 if (max_score2 == score2)
1556 variant2 = NULL_TREE;
1557 else if (score2 > max_score2)
1558 {
1559 max_score2 = score2;
1560 variant2 = defer[i] ? NULL_TREE : attr1;
1561 }
1562 }
1563 }
1564
1565 /* If there is a clear winner variant with the score which is not
1566 deferred, verify it is not a strict subset of any other context
1567 selector and if it is not, it is the best alternative no matter
1568 whether the others do or don't match. */
1569 if (variant1 && variant1 == variant2)
1570 {
1571 tree ctx1 = TREE_VALUE (TREE_VALUE (variant1));
1572 FOR_EACH_VEC_ELT (variants, i, attr2)
1573 {
1574 if (attr2 == variant1)
1575 continue;
1576 tree ctx2 = TREE_VALUE (TREE_VALUE (attr2));
1577 int r = omp_context_selector_compare (ctx1, ctx2);
1578 if (r == -1)
1579 {
1580 /* The winner is a strict subset of ctx2, can't
1581 decide now. */
1582 variant1 = NULL_TREE;
1583 break;
1584 }
1585 }
1586 if (variant1)
1587 return TREE_PURPOSE (TREE_VALUE (variant1));
1588 }
1589
1590 return base;
1591 }
1592
1593 if (variants.length () == 1)
1594 return TREE_PURPOSE (TREE_VALUE (variants[0]));
1595
1596 /* A context selector that is a strict subset of another context selector has a score
1597 of zero. */
1598 tree attr1, attr2;
1599 unsigned int i, j;
1600 FOR_EACH_VEC_ELT (variants, i, attr1)
1601 if (attr1)
1602 {
1603 tree ctx1 = TREE_VALUE (TREE_VALUE (attr1));
1604 FOR_EACH_VEC_ELT_FROM (variants, j, attr2, i + 1)
1605 if (attr2)
1606 {
1607 tree ctx2 = TREE_VALUE (TREE_VALUE (attr2));
1608 int r = omp_context_selector_compare (ctx1, ctx2);
1609 if (r == -1)
1610 {
1611 /* ctx1 is a strict subset of ctx2, remove
1612 attr1 from the vector. */
1613 variants[i] = NULL_TREE;
1614 break;
1615 }
1616 else if (r == 1)
1617 /* ctx2 is a strict subset of ctx1, remove attr2
1618 from the vector. */
1619 variants[j] = NULL_TREE;
1620 }
1621 }
1622 widest_int max_score1 = 0;
1623 widest_int max_score2 = 0;
1624 bool first = true;
1625 FOR_EACH_VEC_ELT (variants, i, attr1)
1626 if (attr1)
1627 {
1628 if (variant1)
1629 {
1630 widest_int score1;
1631 widest_int score2;
1632 bool need_two;
1633 tree ctx;
1634 if (first)
1635 {
1636 first = false;
1637 ctx = TREE_VALUE (TREE_VALUE (variant1));
1638 need_two = omp_context_compute_score (ctx, &max_score1, false);
1639 if (need_two)
1640 omp_context_compute_score (ctx, &max_score2, true);
1641 else
1642 max_score2 = max_score1;
1643 }
1644 ctx = TREE_VALUE (TREE_VALUE (attr1));
1645 need_two = omp_context_compute_score (ctx, &score1, false);
1646 if (need_two)
1647 omp_context_compute_score (ctx, &score2, true);
1648 else
1649 score2 = score1;
1650 if (score1 > max_score1)
1651 {
1652 max_score1 = score1;
1653 variant1 = attr1;
1654 }
1655 if (score2 > max_score2)
1656 {
1657 max_score2 = score2;
1658 variant2 = attr1;
1659 }
1660 }
1661 else
1662 {
1663 variant1 = attr1;
1664 variant2 = attr1;
1665 }
1666 }
1667 /* If there is a disagreement on which variant has the highest score
1668 depending on whether it will be in a declare simd clone or not,
1669 punt for now and defer until after IPA where we will know that. */
1670 return ((variant1 && variant1 == variant2)
1671 ? TREE_PURPOSE (TREE_VALUE (variant1)) : base);
1672 }
1673
1674
1675 /* Encode an oacc launch argument. This matches the GOMP_LAUNCH_PACK
1676 macro on gomp-constants.h. We do not check for overflow. */
1677
1678 tree
oacc_launch_pack(unsigned code,tree device,unsigned op)1679 oacc_launch_pack (unsigned code, tree device, unsigned op)
1680 {
1681 tree res;
1682
1683 res = build_int_cst (unsigned_type_node, GOMP_LAUNCH_PACK (code, 0, op));
1684 if (device)
1685 {
1686 device = fold_build2 (LSHIFT_EXPR, unsigned_type_node,
1687 device, build_int_cst (unsigned_type_node,
1688 GOMP_LAUNCH_DEVICE_SHIFT));
1689 res = fold_build2 (BIT_IOR_EXPR, unsigned_type_node, res, device);
1690 }
1691 return res;
1692 }
1693
1694 /* FIXME: What is the following comment for? */
1695 /* Look for compute grid dimension clauses and convert to an attribute
1696 attached to FN. This permits the target-side code to (a) massage
1697 the dimensions, (b) emit that data and (c) optimize. Non-constant
1698 dimensions are pushed onto ARGS.
1699
1700 The attribute value is a TREE_LIST. A set of dimensions is
1701 represented as a list of INTEGER_CST. Those that are runtime
1702 exprs are represented as an INTEGER_CST of zero.
1703
1704 TODO: Normally the attribute will just contain a single such list. If
1705 however it contains a list of lists, this will represent the use of
1706 device_type. Each member of the outer list is an assoc list of
1707 dimensions, keyed by the device type. The first entry will be the
1708 default. Well, that's the plan. */
1709
1710 /* Replace any existing oacc fn attribute with updated dimensions. */
1711
1712 /* Variant working on a list of attributes. */
1713
1714 tree
oacc_replace_fn_attrib_attr(tree attribs,tree dims)1715 oacc_replace_fn_attrib_attr (tree attribs, tree dims)
1716 {
1717 tree ident = get_identifier (OACC_FN_ATTRIB);
1718
1719 /* If we happen to be present as the first attrib, drop it. */
1720 if (attribs && TREE_PURPOSE (attribs) == ident)
1721 attribs = TREE_CHAIN (attribs);
1722 return tree_cons (ident, dims, attribs);
1723 }
1724
1725 /* Variant working on a function decl. */
1726
1727 void
oacc_replace_fn_attrib(tree fn,tree dims)1728 oacc_replace_fn_attrib (tree fn, tree dims)
1729 {
1730 DECL_ATTRIBUTES (fn)
1731 = oacc_replace_fn_attrib_attr (DECL_ATTRIBUTES (fn), dims);
1732 }
1733
1734 /* Scan CLAUSES for launch dimensions and attach them to the oacc
1735 function attribute. Push any that are non-constant onto the ARGS
1736 list, along with an appropriate GOMP_LAUNCH_DIM tag. */
1737
1738 void
oacc_set_fn_attrib(tree fn,tree clauses,vec<tree> * args)1739 oacc_set_fn_attrib (tree fn, tree clauses, vec<tree> *args)
1740 {
1741 /* Must match GOMP_DIM ordering. */
1742 static const omp_clause_code ids[]
1743 = { OMP_CLAUSE_NUM_GANGS, OMP_CLAUSE_NUM_WORKERS,
1744 OMP_CLAUSE_VECTOR_LENGTH };
1745 unsigned ix;
1746 tree dims[GOMP_DIM_MAX];
1747
1748 tree attr = NULL_TREE;
1749 unsigned non_const = 0;
1750
1751 for (ix = GOMP_DIM_MAX; ix--;)
1752 {
1753 tree clause = omp_find_clause (clauses, ids[ix]);
1754 tree dim = NULL_TREE;
1755
1756 if (clause)
1757 dim = OMP_CLAUSE_EXPR (clause, ids[ix]);
1758 dims[ix] = dim;
1759 if (dim && TREE_CODE (dim) != INTEGER_CST)
1760 {
1761 dim = integer_zero_node;
1762 non_const |= GOMP_DIM_MASK (ix);
1763 }
1764 attr = tree_cons (NULL_TREE, dim, attr);
1765 }
1766
1767 oacc_replace_fn_attrib (fn, attr);
1768
1769 if (non_const)
1770 {
1771 /* Push a dynamic argument set. */
1772 args->safe_push (oacc_launch_pack (GOMP_LAUNCH_DIM,
1773 NULL_TREE, non_const));
1774 for (unsigned ix = 0; ix != GOMP_DIM_MAX; ix++)
1775 if (non_const & GOMP_DIM_MASK (ix))
1776 args->safe_push (dims[ix]);
1777 }
1778 }
1779
1780 /* Verify OpenACC routine clauses.
1781
1782 Returns 0 if FNDECL should be marked with an OpenACC 'routine' directive, 1
1783 if it has already been marked in compatible way, and -1 if incompatible.
1784 Upon returning, the chain of clauses will contain exactly one clause
1785 specifying the level of parallelism. */
1786
1787 int
oacc_verify_routine_clauses(tree fndecl,tree * clauses,location_t loc,const char * routine_str)1788 oacc_verify_routine_clauses (tree fndecl, tree *clauses, location_t loc,
1789 const char *routine_str)
1790 {
1791 tree c_level = NULL_TREE;
1792 tree c_p = NULL_TREE;
1793 for (tree c = *clauses; c; c_p = c, c = OMP_CLAUSE_CHAIN (c))
1794 switch (OMP_CLAUSE_CODE (c))
1795 {
1796 case OMP_CLAUSE_GANG:
1797 case OMP_CLAUSE_WORKER:
1798 case OMP_CLAUSE_VECTOR:
1799 case OMP_CLAUSE_SEQ:
1800 if (c_level == NULL_TREE)
1801 c_level = c;
1802 else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_CODE (c_level))
1803 {
1804 /* This has already been diagnosed in the front ends. */
1805 /* Drop the duplicate clause. */
1806 gcc_checking_assert (c_p != NULL_TREE);
1807 OMP_CLAUSE_CHAIN (c_p) = OMP_CLAUSE_CHAIN (c);
1808 c = c_p;
1809 }
1810 else
1811 {
1812 error_at (OMP_CLAUSE_LOCATION (c),
1813 "%qs specifies a conflicting level of parallelism",
1814 omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
1815 inform (OMP_CLAUSE_LOCATION (c_level),
1816 "... to the previous %qs clause here",
1817 omp_clause_code_name[OMP_CLAUSE_CODE (c_level)]);
1818 /* Drop the conflicting clause. */
1819 gcc_checking_assert (c_p != NULL_TREE);
1820 OMP_CLAUSE_CHAIN (c_p) = OMP_CLAUSE_CHAIN (c);
1821 c = c_p;
1822 }
1823 break;
1824 default:
1825 gcc_unreachable ();
1826 }
1827 if (c_level == NULL_TREE)
1828 {
1829 /* Default to an implicit 'seq' clause. */
1830 c_level = build_omp_clause (loc, OMP_CLAUSE_SEQ);
1831 OMP_CLAUSE_CHAIN (c_level) = *clauses;
1832 *clauses = c_level;
1833 }
1834 /* In *clauses, we now have exactly one clause specifying the level of
1835 parallelism. */
1836
1837 tree attr
1838 = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (fndecl));
1839 if (attr != NULL_TREE)
1840 {
1841 /* Diagnose if "#pragma omp declare target" has also been applied. */
1842 if (TREE_VALUE (attr) == NULL_TREE)
1843 {
1844 /* See <https://gcc.gnu.org/PR93465>; the semantics of combining
1845 OpenACC and OpenMP 'target' are not clear. */
1846 error_at (loc,
1847 "cannot apply %<%s%> to %qD, which has also been"
1848 " marked with an OpenMP 'declare target' directive",
1849 routine_str, fndecl);
1850 /* Incompatible. */
1851 return -1;
1852 }
1853
1854 /* If a "#pragma acc routine" has already been applied, just verify
1855 this one for compatibility. */
1856 /* Collect previous directive's clauses. */
1857 tree c_level_p = NULL_TREE;
1858 for (tree c = TREE_VALUE (attr); c; c = OMP_CLAUSE_CHAIN (c))
1859 switch (OMP_CLAUSE_CODE (c))
1860 {
1861 case OMP_CLAUSE_GANG:
1862 case OMP_CLAUSE_WORKER:
1863 case OMP_CLAUSE_VECTOR:
1864 case OMP_CLAUSE_SEQ:
1865 gcc_checking_assert (c_level_p == NULL_TREE);
1866 c_level_p = c;
1867 break;
1868 default:
1869 gcc_unreachable ();
1870 }
1871 gcc_checking_assert (c_level_p != NULL_TREE);
1872 /* ..., and compare to current directive's, which we've already collected
1873 above. */
1874 tree c_diag;
1875 tree c_diag_p;
1876 /* Matching level of parallelism? */
1877 if (OMP_CLAUSE_CODE (c_level) != OMP_CLAUSE_CODE (c_level_p))
1878 {
1879 c_diag = c_level;
1880 c_diag_p = c_level_p;
1881 goto incompatible;
1882 }
1883 /* Compatible. */
1884 return 1;
1885
1886 incompatible:
1887 if (c_diag != NULL_TREE)
1888 error_at (OMP_CLAUSE_LOCATION (c_diag),
1889 "incompatible %qs clause when applying"
1890 " %<%s%> to %qD, which has already been"
1891 " marked with an OpenACC 'routine' directive",
1892 omp_clause_code_name[OMP_CLAUSE_CODE (c_diag)],
1893 routine_str, fndecl);
1894 else if (c_diag_p != NULL_TREE)
1895 error_at (loc,
1896 "missing %qs clause when applying"
1897 " %<%s%> to %qD, which has already been"
1898 " marked with an OpenACC 'routine' directive",
1899 omp_clause_code_name[OMP_CLAUSE_CODE (c_diag_p)],
1900 routine_str, fndecl);
1901 else
1902 gcc_unreachable ();
1903 if (c_diag_p != NULL_TREE)
1904 inform (OMP_CLAUSE_LOCATION (c_diag_p),
1905 "... with %qs clause here",
1906 omp_clause_code_name[OMP_CLAUSE_CODE (c_diag_p)]);
1907 else
1908 {
1909 /* In the front ends, we don't preserve location information for the
1910 OpenACC routine directive itself. However, that of c_level_p
1911 should be close. */
1912 location_t loc_routine = OMP_CLAUSE_LOCATION (c_level_p);
1913 inform (loc_routine, "... without %qs clause near to here",
1914 omp_clause_code_name[OMP_CLAUSE_CODE (c_diag)]);
1915 }
1916 /* Incompatible. */
1917 return -1;
1918 }
1919
1920 return 0;
1921 }
1922
1923 /* Process the OpenACC 'routine' directive clauses to generate an attribute
1924 for the level of parallelism. All dimensions have a size of zero
1925 (dynamic). TREE_PURPOSE is set to indicate whether that dimension
1926 can have a loop partitioned on it. non-zero indicates
1927 yes, zero indicates no. By construction once a non-zero has been
1928 reached, further inner dimensions must also be non-zero. We set
1929 TREE_VALUE to zero for the dimensions that may be partitioned and
1930 1 for the other ones -- if a loop is (erroneously) spawned at
1931 an outer level, we don't want to try and partition it. */
1932
1933 tree
oacc_build_routine_dims(tree clauses)1934 oacc_build_routine_dims (tree clauses)
1935 {
1936 /* Must match GOMP_DIM ordering. */
1937 static const omp_clause_code ids[]
1938 = {OMP_CLAUSE_GANG, OMP_CLAUSE_WORKER, OMP_CLAUSE_VECTOR, OMP_CLAUSE_SEQ};
1939 int ix;
1940 int level = -1;
1941
1942 for (; clauses; clauses = OMP_CLAUSE_CHAIN (clauses))
1943 for (ix = GOMP_DIM_MAX + 1; ix--;)
1944 if (OMP_CLAUSE_CODE (clauses) == ids[ix])
1945 {
1946 level = ix;
1947 break;
1948 }
1949 gcc_checking_assert (level >= 0);
1950
1951 tree dims = NULL_TREE;
1952
1953 for (ix = GOMP_DIM_MAX; ix--;)
1954 dims = tree_cons (build_int_cst (boolean_type_node, ix >= level),
1955 build_int_cst (integer_type_node, ix < level), dims);
1956
1957 return dims;
1958 }
1959
1960 /* Retrieve the oacc function attrib and return it. Non-oacc
1961 functions will return NULL. */
1962
1963 tree
oacc_get_fn_attrib(tree fn)1964 oacc_get_fn_attrib (tree fn)
1965 {
1966 return lookup_attribute (OACC_FN_ATTRIB, DECL_ATTRIBUTES (fn));
1967 }
1968
1969 /* Return true if FN is an OpenMP or OpenACC offloading function. */
1970
1971 bool
offloading_function_p(tree fn)1972 offloading_function_p (tree fn)
1973 {
1974 tree attrs = DECL_ATTRIBUTES (fn);
1975 return (lookup_attribute ("omp declare target", attrs)
1976 || lookup_attribute ("omp target entrypoint", attrs));
1977 }
1978
1979 /* Extract an oacc execution dimension from FN. FN must be an
1980 offloaded function or routine that has already had its execution
1981 dimensions lowered to the target-specific values. */
1982
1983 int
oacc_get_fn_dim_size(tree fn,int axis)1984 oacc_get_fn_dim_size (tree fn, int axis)
1985 {
1986 tree attrs = oacc_get_fn_attrib (fn);
1987
1988 gcc_assert (axis < GOMP_DIM_MAX);
1989
1990 tree dims = TREE_VALUE (attrs);
1991 while (axis--)
1992 dims = TREE_CHAIN (dims);
1993
1994 int size = TREE_INT_CST_LOW (TREE_VALUE (dims));
1995
1996 return size;
1997 }
1998
1999 /* Extract the dimension axis from an IFN_GOACC_DIM_POS or
2000 IFN_GOACC_DIM_SIZE call. */
2001
2002 int
oacc_get_ifn_dim_arg(const gimple * stmt)2003 oacc_get_ifn_dim_arg (const gimple *stmt)
2004 {
2005 gcc_checking_assert (gimple_call_internal_fn (stmt) == IFN_GOACC_DIM_SIZE
2006 || gimple_call_internal_fn (stmt) == IFN_GOACC_DIM_POS);
2007 tree arg = gimple_call_arg (stmt, 0);
2008 HOST_WIDE_INT axis = TREE_INT_CST_LOW (arg);
2009
2010 gcc_checking_assert (axis >= 0 && axis < GOMP_DIM_MAX);
2011 return (int) axis;
2012 }
2013