1 /* General types and functions that are uselful for processing of OpenMP,
2 OpenACC and similar directivers at various stages of compilation.
4 Copyright (C) 2005-2021 Free Software Foundation, Inc.
6 This file is part of GCC.
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
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
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/>. */
22 /* Find an OMP clause of type KIND within CLAUSES. */
26 #include "coretypes.h"
32 #include "diagnostic-core.h"
33 #include "fold-const.h"
34 #include "langhooks.h"
35 #include "omp-general.h"
36 #include "stringpool.h"
40 #include "alloc-pool.h"
41 #include "symbol-summary.h"
42 #include "tree-pass.h"
43 #include "omp-device-properties.h"
44 #include "tree-iterator.h"
45 #include "data-streamer.h"
46 #include "streamer-hooks.h"
49 enum omp_requires omp_requires_mask
;
52 omp_find_clause (tree clauses
, enum omp_clause_code kind
)
54 for (; clauses
; clauses
= OMP_CLAUSE_CHAIN (clauses
))
55 if (OMP_CLAUSE_CODE (clauses
) == kind
)
61 /* True if OpenMP should regard this DECL as being a scalar which has Fortran's
62 allocatable or pointer attribute. */
64 omp_is_allocatable_or_ptr (tree decl
)
66 return lang_hooks
.decls
.omp_is_allocatable_or_ptr (decl
);
69 /* Check whether this DECL belongs to a Fortran optional argument.
70 With 'for_present_check' set to false, decls which are optional parameters
71 themselve are returned as tree - or a NULL_TREE otherwise. Those decls are
72 always pointers. With 'for_present_check' set to true, the decl for checking
73 whether an argument is present is returned; for arguments with value
74 attribute this is the hidden argument and of BOOLEAN_TYPE. If the decl is
75 unrelated to optional arguments, NULL_TREE is returned. */
78 omp_check_optional_argument (tree decl
, bool for_present_check
)
80 return lang_hooks
.decls
.omp_check_optional_argument (decl
, for_present_check
);
83 /* True if OpenMP should privatize what this DECL points to rather
84 than the DECL itself. */
87 omp_privatize_by_reference (tree decl
)
89 return lang_hooks
.decls
.omp_privatize_by_reference (decl
);
92 /* Adjust *COND_CODE and *N2 so that the former is either LT_EXPR or GT_EXPR,
93 given that V is the loop index variable and STEP is loop step. */
96 omp_adjust_for_condition (location_t loc
, enum tree_code
*cond_code
, tree
*n2
,
106 gcc_assert (TREE_CODE (step
) == INTEGER_CST
);
107 if (TREE_CODE (TREE_TYPE (v
)) == INTEGER_TYPE
)
109 if (integer_onep (step
))
110 *cond_code
= LT_EXPR
;
113 gcc_assert (integer_minus_onep (step
));
114 *cond_code
= GT_EXPR
;
119 tree unit
= TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (v
)));
120 gcc_assert (TREE_CODE (unit
) == INTEGER_CST
);
121 if (tree_int_cst_equal (unit
, step
))
122 *cond_code
= LT_EXPR
;
125 gcc_assert (wi::neg (wi::to_widest (unit
))
126 == wi::to_widest (step
));
127 *cond_code
= GT_EXPR
;
134 if (POINTER_TYPE_P (TREE_TYPE (*n2
)))
135 *n2
= fold_build_pointer_plus_hwi_loc (loc
, *n2
, 1);
137 *n2
= fold_build2_loc (loc
, PLUS_EXPR
, TREE_TYPE (*n2
), *n2
,
138 build_int_cst (TREE_TYPE (*n2
), 1));
139 *cond_code
= LT_EXPR
;
142 if (POINTER_TYPE_P (TREE_TYPE (*n2
)))
143 *n2
= fold_build_pointer_plus_hwi_loc (loc
, *n2
, -1);
145 *n2
= fold_build2_loc (loc
, MINUS_EXPR
, TREE_TYPE (*n2
), *n2
,
146 build_int_cst (TREE_TYPE (*n2
), 1));
147 *cond_code
= GT_EXPR
;
154 /* Return the looping step from INCR, extracted from the step of a gimple omp
158 omp_get_for_step_from_incr (location_t loc
, tree incr
)
161 switch (TREE_CODE (incr
))
164 step
= TREE_OPERAND (incr
, 1);
166 case POINTER_PLUS_EXPR
:
167 step
= fold_convert (ssizetype
, TREE_OPERAND (incr
, 1));
170 step
= TREE_OPERAND (incr
, 1);
171 step
= fold_build1_loc (loc
, NEGATE_EXPR
, TREE_TYPE (step
), step
);
179 /* Extract the header elements of parallel loop FOR_STMT and store
183 omp_extract_for_data (gomp_for
*for_stmt
, struct omp_for_data
*fd
,
184 struct omp_for_data_loop
*loops
)
186 tree t
, var
, *collapse_iter
, *collapse_count
;
187 tree count
= NULL_TREE
, iter_type
= long_integer_type_node
;
188 struct omp_for_data_loop
*loop
;
190 struct omp_for_data_loop dummy_loop
;
191 location_t loc
= gimple_location (for_stmt
);
192 bool simd
= gimple_omp_for_kind (for_stmt
) == GF_OMP_FOR_KIND_SIMD
;
193 bool distribute
= gimple_omp_for_kind (for_stmt
)
194 == GF_OMP_FOR_KIND_DISTRIBUTE
;
195 bool taskloop
= gimple_omp_for_kind (for_stmt
)
196 == GF_OMP_FOR_KIND_TASKLOOP
;
197 bool order_reproducible
= false;
200 fd
->for_stmt
= for_stmt
;
202 fd
->have_nowait
= distribute
|| simd
;
203 fd
->have_ordered
= false;
204 fd
->have_reductemp
= false;
205 fd
->have_pointer_condtemp
= false;
206 fd
->have_scantemp
= false;
207 fd
->have_nonctrl_scantemp
= false;
208 fd
->non_rect
= false;
209 fd
->lastprivate_conditional
= 0;
210 fd
->tiling
= NULL_TREE
;
213 fd
->first_nonrect
= -1;
214 fd
->last_nonrect
= -1;
215 fd
->sched_kind
= OMP_CLAUSE_SCHEDULE_STATIC
;
216 fd
->sched_modifiers
= 0;
217 fd
->chunk_size
= NULL_TREE
;
218 fd
->simd_schedule
= false;
219 fd
->first_inner_iterations
= NULL_TREE
;
220 fd
->factor
= NULL_TREE
;
221 fd
->adjn1
= NULL_TREE
;
222 collapse_iter
= NULL
;
223 collapse_count
= NULL
;
225 for (t
= gimple_omp_for_clauses (for_stmt
); t
; t
= OMP_CLAUSE_CHAIN (t
))
226 switch (OMP_CLAUSE_CODE (t
))
228 case OMP_CLAUSE_NOWAIT
:
229 fd
->have_nowait
= true;
231 case OMP_CLAUSE_ORDERED
:
232 fd
->have_ordered
= true;
233 if (OMP_CLAUSE_ORDERED_EXPR (t
))
234 fd
->ordered
= tree_to_shwi (OMP_CLAUSE_ORDERED_EXPR (t
));
236 case OMP_CLAUSE_SCHEDULE
:
237 gcc_assert (!distribute
&& !taskloop
);
239 = (enum omp_clause_schedule_kind
)
240 (OMP_CLAUSE_SCHEDULE_KIND (t
) & OMP_CLAUSE_SCHEDULE_MASK
);
241 fd
->sched_modifiers
= (OMP_CLAUSE_SCHEDULE_KIND (t
)
242 & ~OMP_CLAUSE_SCHEDULE_MASK
);
243 fd
->chunk_size
= OMP_CLAUSE_SCHEDULE_CHUNK_EXPR (t
);
244 fd
->simd_schedule
= OMP_CLAUSE_SCHEDULE_SIMD (t
);
246 case OMP_CLAUSE_DIST_SCHEDULE
:
247 gcc_assert (distribute
);
248 fd
->chunk_size
= OMP_CLAUSE_DIST_SCHEDULE_CHUNK_EXPR (t
);
250 case OMP_CLAUSE_COLLAPSE
:
251 fd
->collapse
= tree_to_shwi (OMP_CLAUSE_COLLAPSE_EXPR (t
));
252 if (fd
->collapse
> 1)
254 collapse_iter
= &OMP_CLAUSE_COLLAPSE_ITERVAR (t
);
255 collapse_count
= &OMP_CLAUSE_COLLAPSE_COUNT (t
);
258 case OMP_CLAUSE_TILE
:
259 fd
->tiling
= OMP_CLAUSE_TILE_LIST (t
);
260 fd
->collapse
= list_length (fd
->tiling
);
261 gcc_assert (fd
->collapse
);
262 collapse_iter
= &OMP_CLAUSE_TILE_ITERVAR (t
);
263 collapse_count
= &OMP_CLAUSE_TILE_COUNT (t
);
265 case OMP_CLAUSE__REDUCTEMP_
:
266 fd
->have_reductemp
= true;
268 case OMP_CLAUSE_LASTPRIVATE
:
269 if (OMP_CLAUSE_LASTPRIVATE_CONDITIONAL (t
))
270 fd
->lastprivate_conditional
++;
272 case OMP_CLAUSE__CONDTEMP_
:
273 if (POINTER_TYPE_P (TREE_TYPE (OMP_CLAUSE_DECL (t
))))
274 fd
->have_pointer_condtemp
= true;
276 case OMP_CLAUSE__SCANTEMP_
:
277 fd
->have_scantemp
= true;
278 if (!OMP_CLAUSE__SCANTEMP__ALLOC (t
)
279 && !OMP_CLAUSE__SCANTEMP__CONTROL (t
))
280 fd
->have_nonctrl_scantemp
= true;
282 case OMP_CLAUSE_ORDER
:
283 /* FIXME: For OpenMP 5.2 this should change to
284 if (OMP_CLAUSE_ORDER_REPRODUCIBLE (t))
285 (with the exception of loop construct but that lowers to
286 no schedule/dist_schedule clauses currently). */
287 if (!OMP_CLAUSE_ORDER_UNCONSTRAINED (t
))
288 order_reproducible
= true;
293 /* For order(reproducible:concurrent) schedule ({dynamic,guided,runtime})
294 we have either the option to expensively remember at runtime how we've
295 distributed work from first loop and reuse that in following loops with
296 the same number of iterations and schedule, or just force static schedule.
297 OpenMP API calls etc. aren't allowed in order(concurrent) bodies so
298 users can't observe it easily anyway. */
299 if (order_reproducible
)
300 fd
->sched_kind
= OMP_CLAUSE_SCHEDULE_STATIC
;
301 if (fd
->collapse
> 1 || fd
->tiling
)
304 fd
->loops
= &fd
->loop
;
306 if (fd
->ordered
&& fd
->collapse
== 1 && loops
!= NULL
)
311 collapse_iter
= &iterv
;
312 collapse_count
= &countv
;
315 /* FIXME: for now map schedule(auto) to schedule(static).
316 There should be analysis to determine whether all iterations
317 are approximately the same amount of work (then schedule(static)
318 is best) or if it varies (then schedule(dynamic,N) is better). */
319 if (fd
->sched_kind
== OMP_CLAUSE_SCHEDULE_AUTO
)
321 fd
->sched_kind
= OMP_CLAUSE_SCHEDULE_STATIC
;
322 gcc_assert (fd
->chunk_size
== NULL
);
324 gcc_assert ((fd
->collapse
== 1 && !fd
->tiling
) || collapse_iter
!= NULL
);
326 fd
->sched_kind
= OMP_CLAUSE_SCHEDULE_RUNTIME
;
327 if (fd
->sched_kind
== OMP_CLAUSE_SCHEDULE_RUNTIME
)
328 gcc_assert (fd
->chunk_size
== NULL
);
329 else if (fd
->chunk_size
== NULL
)
331 /* We only need to compute a default chunk size for ordered
332 static loops and dynamic loops. */
333 if (fd
->sched_kind
!= OMP_CLAUSE_SCHEDULE_STATIC
335 fd
->chunk_size
= (fd
->sched_kind
== OMP_CLAUSE_SCHEDULE_STATIC
)
336 ? integer_zero_node
: integer_one_node
;
339 int cnt
= fd
->ordered
? fd
->ordered
: fd
->collapse
;
340 int single_nonrect
= -1;
341 tree single_nonrect_count
= NULL_TREE
;
342 enum tree_code single_nonrect_cond_code
= ERROR_MARK
;
343 for (i
= 1; i
< cnt
; i
++)
345 tree n1
= gimple_omp_for_initial (for_stmt
, i
);
346 tree n2
= gimple_omp_for_final (for_stmt
, i
);
347 if (TREE_CODE (n1
) == TREE_VEC
)
354 for (int j
= i
- 1; j
>= 0; j
--)
355 if (TREE_VEC_ELT (n1
, 0) == gimple_omp_for_index (for_stmt
, j
))
362 else if (TREE_CODE (n2
) == TREE_VEC
)
369 for (int j
= i
- 1; j
>= 0; j
--)
370 if (TREE_VEC_ELT (n2
, 0) == gimple_omp_for_index (for_stmt
, j
))
378 for (i
= 0; i
< cnt
; i
++)
383 && (fd
->ordered
== 0 || loops
== NULL
))
385 else if (loops
!= NULL
)
390 loop
->v
= gimple_omp_for_index (for_stmt
, i
);
391 gcc_assert (SSA_VAR_P (loop
->v
));
392 gcc_assert (TREE_CODE (TREE_TYPE (loop
->v
)) == INTEGER_TYPE
393 || TREE_CODE (TREE_TYPE (loop
->v
)) == POINTER_TYPE
);
394 var
= TREE_CODE (loop
->v
) == SSA_NAME
? SSA_NAME_VAR (loop
->v
) : loop
->v
;
395 loop
->n1
= gimple_omp_for_initial (for_stmt
, i
);
396 loop
->m1
= NULL_TREE
;
397 loop
->m2
= NULL_TREE
;
399 loop
->non_rect_referenced
= false;
400 if (TREE_CODE (loop
->n1
) == TREE_VEC
)
402 for (int j
= i
- 1; j
>= 0; j
--)
403 if (TREE_VEC_ELT (loop
->n1
, 0) == gimple_omp_for_index (for_stmt
, j
))
407 loops
[j
].non_rect_referenced
= true;
408 if (fd
->first_nonrect
== -1 || fd
->first_nonrect
> j
)
409 fd
->first_nonrect
= j
;
412 gcc_assert (loop
->outer
);
413 loop
->m1
= TREE_VEC_ELT (loop
->n1
, 1);
414 loop
->n1
= TREE_VEC_ELT (loop
->n1
, 2);
416 fd
->last_nonrect
= i
;
419 loop
->cond_code
= gimple_omp_for_cond (for_stmt
, i
);
420 loop
->n2
= gimple_omp_for_final (for_stmt
, i
);
421 gcc_assert (loop
->cond_code
!= NE_EXPR
422 || (gimple_omp_for_kind (for_stmt
)
423 != GF_OMP_FOR_KIND_OACC_LOOP
));
424 if (TREE_CODE (loop
->n2
) == TREE_VEC
)
427 gcc_assert (TREE_VEC_ELT (loop
->n2
, 0)
428 == gimple_omp_for_index (for_stmt
, i
- loop
->outer
));
430 for (int j
= i
- 1; j
>= 0; j
--)
431 if (TREE_VEC_ELT (loop
->n2
, 0) == gimple_omp_for_index (for_stmt
, j
))
435 loops
[j
].non_rect_referenced
= true;
436 if (fd
->first_nonrect
== -1 || fd
->first_nonrect
> j
)
437 fd
->first_nonrect
= j
;
440 gcc_assert (loop
->outer
);
441 loop
->m2
= TREE_VEC_ELT (loop
->n2
, 1);
442 loop
->n2
= TREE_VEC_ELT (loop
->n2
, 2);
444 fd
->last_nonrect
= i
;
447 t
= gimple_omp_for_incr (for_stmt
, i
);
448 gcc_assert (TREE_OPERAND (t
, 0) == var
);
449 loop
->step
= omp_get_for_step_from_incr (loc
, t
);
451 omp_adjust_for_condition (loc
, &loop
->cond_code
, &loop
->n2
, loop
->v
,
455 || (fd
->sched_kind
== OMP_CLAUSE_SCHEDULE_STATIC
456 && !fd
->have_ordered
))
458 if (fd
->collapse
== 1 && !fd
->tiling
)
459 iter_type
= TREE_TYPE (loop
->v
);
461 || TYPE_PRECISION (iter_type
)
462 < TYPE_PRECISION (TREE_TYPE (loop
->v
)))
464 = build_nonstandard_integer_type
465 (TYPE_PRECISION (TREE_TYPE (loop
->v
)), 1);
467 else if (iter_type
!= long_long_unsigned_type_node
)
469 if (POINTER_TYPE_P (TREE_TYPE (loop
->v
)))
470 iter_type
= long_long_unsigned_type_node
;
471 else if (TYPE_UNSIGNED (TREE_TYPE (loop
->v
))
472 && TYPE_PRECISION (TREE_TYPE (loop
->v
))
473 >= TYPE_PRECISION (iter_type
))
477 if (loop
->cond_code
== LT_EXPR
)
478 n
= fold_build2_loc (loc
, PLUS_EXPR
, TREE_TYPE (loop
->v
),
479 loop
->n2
, loop
->step
);
484 || TREE_CODE (n
) != INTEGER_CST
485 || tree_int_cst_lt (TYPE_MAX_VALUE (iter_type
), n
))
486 iter_type
= long_long_unsigned_type_node
;
488 else if (TYPE_PRECISION (TREE_TYPE (loop
->v
))
489 > TYPE_PRECISION (iter_type
))
493 if (loop
->cond_code
== LT_EXPR
)
496 n2
= fold_build2_loc (loc
, PLUS_EXPR
, TREE_TYPE (loop
->v
),
497 loop
->n2
, loop
->step
);
501 n1
= fold_build2_loc (loc
, MINUS_EXPR
, TREE_TYPE (loop
->v
),
502 loop
->n2
, loop
->step
);
507 || TREE_CODE (n1
) != INTEGER_CST
508 || TREE_CODE (n2
) != INTEGER_CST
509 || !tree_int_cst_lt (TYPE_MIN_VALUE (iter_type
), n1
)
510 || !tree_int_cst_lt (n2
, TYPE_MAX_VALUE (iter_type
)))
511 iter_type
= long_long_unsigned_type_node
;
515 if (i
>= fd
->collapse
)
518 if (collapse_count
&& *collapse_count
== NULL
)
520 if (count
&& integer_zerop (count
))
522 tree n1first
= NULL_TREE
, n2first
= NULL_TREE
;
523 tree n1last
= NULL_TREE
, n2last
= NULL_TREE
;
524 tree ostep
= NULL_TREE
;
525 if (loop
->m1
|| loop
->m2
)
527 if (count
== NULL_TREE
)
529 if (single_nonrect
== -1
530 || (loop
->m1
&& TREE_CODE (loop
->m1
) != INTEGER_CST
)
531 || (loop
->m2
&& TREE_CODE (loop
->m2
) != INTEGER_CST
)
532 || TREE_CODE (loop
->n1
) != INTEGER_CST
533 || TREE_CODE (loop
->n2
) != INTEGER_CST
534 || TREE_CODE (loop
->step
) != INTEGER_CST
)
539 tree var
= gimple_omp_for_initial (for_stmt
, single_nonrect
);
540 tree itype
= TREE_TYPE (var
);
541 tree first
= gimple_omp_for_initial (for_stmt
, single_nonrect
);
542 t
= gimple_omp_for_incr (for_stmt
, single_nonrect
);
543 ostep
= omp_get_for_step_from_incr (loc
, t
);
544 t
= fold_binary (MINUS_EXPR
, long_long_unsigned_type_node
,
545 single_nonrect_count
,
546 build_one_cst (long_long_unsigned_type_node
));
547 t
= fold_convert (itype
, t
);
548 first
= fold_convert (itype
, first
);
549 ostep
= fold_convert (itype
, ostep
);
550 tree last
= fold_binary (PLUS_EXPR
, itype
, first
,
551 fold_binary (MULT_EXPR
, itype
, t
,
553 if (TREE_CODE (first
) != INTEGER_CST
554 || TREE_CODE (last
) != INTEGER_CST
)
561 tree m1
= fold_convert (itype
, loop
->m1
);
562 tree n1
= fold_convert (itype
, loop
->n1
);
563 n1first
= fold_binary (PLUS_EXPR
, itype
,
564 fold_binary (MULT_EXPR
, itype
,
566 n1last
= fold_binary (PLUS_EXPR
, itype
,
567 fold_binary (MULT_EXPR
, itype
,
571 n1first
= n1last
= loop
->n1
;
574 tree n2
= fold_convert (itype
, loop
->n2
);
575 tree m2
= fold_convert (itype
, loop
->m2
);
576 n2first
= fold_binary (PLUS_EXPR
, itype
,
577 fold_binary (MULT_EXPR
, itype
,
579 n2last
= fold_binary (PLUS_EXPR
, itype
,
580 fold_binary (MULT_EXPR
, itype
,
584 n2first
= n2last
= loop
->n2
;
585 n1first
= fold_convert (TREE_TYPE (loop
->v
), n1first
);
586 n2first
= fold_convert (TREE_TYPE (loop
->v
), n2first
);
587 n1last
= fold_convert (TREE_TYPE (loop
->v
), n1last
);
588 n2last
= fold_convert (TREE_TYPE (loop
->v
), n2last
);
589 t
= fold_binary (loop
->cond_code
, boolean_type_node
,
591 tree t2
= fold_binary (loop
->cond_code
, boolean_type_node
,
593 if (t
&& t2
&& integer_nonzerop (t
) && integer_nonzerop (t2
))
594 /* All outer loop iterators have at least one inner loop
595 iteration. Try to compute the count at compile time. */
597 else if (t
&& t2
&& integer_zerop (t
) && integer_zerop (t2
))
598 /* No iterations of the inner loop. count will be set to
600 else if (TYPE_UNSIGNED (itype
)
603 || TREE_CODE (t
) != INTEGER_CST
604 || TREE_CODE (t2
) != INTEGER_CST
)
606 /* Punt (for now). */
612 /* Some iterations of the outer loop have zero iterations
613 of the inner loop, while others have at least one.
614 In this case, we need to adjust one of those outer
615 loop bounds. If ADJ_FIRST, we need to adjust outer n1
616 (first), otherwise outer n2 (last). */
617 bool adj_first
= integer_zerop (t
);
618 tree n1
= fold_convert (itype
, loop
->n1
);
619 tree n2
= fold_convert (itype
, loop
->n2
);
620 tree m1
= loop
->m1
? fold_convert (itype
, loop
->m1
)
621 : build_zero_cst (itype
);
622 tree m2
= loop
->m2
? fold_convert (itype
, loop
->m2
)
623 : build_zero_cst (itype
);
624 t
= fold_binary (MINUS_EXPR
, itype
, n1
, n2
);
625 t2
= fold_binary (MINUS_EXPR
, itype
, m2
, m1
);
626 t
= fold_binary (TRUNC_DIV_EXPR
, itype
, t
, t2
);
627 t2
= fold_binary (MINUS_EXPR
, itype
, t
, first
);
628 t2
= fold_binary (TRUNC_MOD_EXPR
, itype
, t2
, ostep
);
629 t
= fold_binary (MINUS_EXPR
, itype
, t
, t2
);
631 = fold_binary (PLUS_EXPR
, itype
, n1
,
632 fold_binary (MULT_EXPR
, itype
, m1
, t
));
634 = fold_binary (PLUS_EXPR
, itype
, n2
,
635 fold_binary (MULT_EXPR
, itype
, m2
, t
));
636 t2
= fold_binary (loop
->cond_code
, boolean_type_node
,
638 tree t3
= fold_binary (MULT_EXPR
, itype
, m1
, ostep
);
639 tree t4
= fold_binary (MULT_EXPR
, itype
, m2
, ostep
);
644 if (integer_nonzerop (t2
))
651 t3
= fold_binary (MINUS_EXPR
, itype
, n1cur
, t3
);
652 t4
= fold_binary (MINUS_EXPR
, itype
, n2cur
, t4
);
653 t3
= fold_binary (loop
->cond_code
,
654 boolean_type_node
, t3
, t4
);
655 gcc_assert (integer_zerop (t3
));
660 t3
= fold_binary (PLUS_EXPR
, itype
, n1cur
, t3
);
661 t4
= fold_binary (PLUS_EXPR
, itype
, n2cur
, t4
);
662 new_first
= fold_binary (PLUS_EXPR
, itype
, t
, ostep
);
667 t3
= fold_binary (loop
->cond_code
,
668 boolean_type_node
, t3
, t4
);
669 gcc_assert (integer_nonzerop (t3
));
672 diff
= fold_binary (MINUS_EXPR
, itype
, new_first
, first
);
679 if (integer_zerop (t2
))
681 t3
= fold_binary (MINUS_EXPR
, itype
, n1cur
, t3
);
682 t4
= fold_binary (MINUS_EXPR
, itype
, n2cur
, t4
);
683 new_last
= fold_binary (MINUS_EXPR
, itype
, t
, ostep
);
688 t3
= fold_binary (loop
->cond_code
,
689 boolean_type_node
, t3
, t4
);
690 gcc_assert (integer_nonzerop (t3
));
700 t3
= fold_binary (PLUS_EXPR
, itype
, n1cur
, t3
);
701 t4
= fold_binary (PLUS_EXPR
, itype
, n2cur
, t4
);
702 t3
= fold_binary (loop
->cond_code
,
703 boolean_type_node
, t3
, t4
);
704 gcc_assert (integer_zerop (t3
));
707 diff
= fold_binary (MINUS_EXPR
, itype
, last
, new_last
);
709 if (TYPE_UNSIGNED (itype
)
710 && single_nonrect_cond_code
== GT_EXPR
)
711 diff
= fold_binary (TRUNC_DIV_EXPR
, itype
,
712 fold_unary (NEGATE_EXPR
, itype
, diff
),
713 fold_unary (NEGATE_EXPR
, itype
,
716 diff
= fold_binary (TRUNC_DIV_EXPR
, itype
, diff
, ostep
);
717 diff
= fold_convert (long_long_unsigned_type_node
, diff
);
719 = fold_binary (MINUS_EXPR
, long_long_unsigned_type_node
,
720 single_nonrect_count
, diff
);
725 t
= fold_binary (loop
->cond_code
, boolean_type_node
,
726 fold_convert (TREE_TYPE (loop
->v
), loop
->n1
),
727 fold_convert (TREE_TYPE (loop
->v
), loop
->n2
));
728 if (t
&& integer_zerop (t
))
729 count
= build_zero_cst (long_long_unsigned_type_node
);
730 else if ((i
== 0 || count
!= NULL_TREE
)
731 && TREE_CODE (TREE_TYPE (loop
->v
)) == INTEGER_TYPE
732 && TREE_CONSTANT (loop
->n1
)
733 && TREE_CONSTANT (loop
->n2
)
734 && TREE_CODE (loop
->step
) == INTEGER_CST
)
736 tree itype
= TREE_TYPE (loop
->v
);
738 if (POINTER_TYPE_P (itype
))
739 itype
= signed_type_for (itype
);
740 t
= build_int_cst (itype
, (loop
->cond_code
== LT_EXPR
? -1 : 1));
741 t
= fold_build2 (PLUS_EXPR
, itype
,
742 fold_convert (itype
, loop
->step
), t
);
745 if (loop
->m1
|| loop
->m2
)
747 gcc_assert (single_nonrect
!= -1);
751 t
= fold_build2 (PLUS_EXPR
, itype
, t
, fold_convert (itype
, n2
));
752 t
= fold_build2 (MINUS_EXPR
, itype
, t
, fold_convert (itype
, n1
));
753 tree step
= fold_convert_loc (loc
, itype
, loop
->step
);
754 if (TYPE_UNSIGNED (itype
) && loop
->cond_code
== GT_EXPR
)
755 t
= fold_build2 (TRUNC_DIV_EXPR
, itype
,
756 fold_build1 (NEGATE_EXPR
, itype
, t
),
757 fold_build1 (NEGATE_EXPR
, itype
, step
));
759 t
= fold_build2 (TRUNC_DIV_EXPR
, itype
, t
, step
);
760 tree llutype
= long_long_unsigned_type_node
;
761 t
= fold_convert (llutype
, t
);
762 if (loop
->m1
|| loop
->m2
)
764 /* t is number of iterations of inner loop at either first
765 or last value of the outer iterator (the one with fewer
767 Compute t2 = ((m2 - m1) * ostep) / step
768 and niters = outer_count * t
769 + t2 * ((outer_count - 1) * outer_count / 2)
771 tree m1
= loop
->m1
? loop
->m1
: integer_zero_node
;
772 tree m2
= loop
->m2
? loop
->m2
: integer_zero_node
;
773 m1
= fold_convert (itype
, m1
);
774 m2
= fold_convert (itype
, m2
);
775 tree t2
= fold_build2 (MINUS_EXPR
, itype
, m2
, m1
);
776 t2
= fold_build2 (MULT_EXPR
, itype
, t2
, ostep
);
777 if (TYPE_UNSIGNED (itype
) && loop
->cond_code
== GT_EXPR
)
778 t2
= fold_build2 (TRUNC_DIV_EXPR
, itype
,
779 fold_build1 (NEGATE_EXPR
, itype
, t2
),
780 fold_build1 (NEGATE_EXPR
, itype
, step
));
782 t2
= fold_build2 (TRUNC_DIV_EXPR
, itype
, t2
, step
);
783 t2
= fold_convert (llutype
, t2
);
784 fd
->first_inner_iterations
= t
;
786 t
= fold_build2 (MULT_EXPR
, llutype
, t
,
787 single_nonrect_count
);
788 tree t3
= fold_build2 (MINUS_EXPR
, llutype
,
789 single_nonrect_count
,
790 build_one_cst (llutype
));
791 t3
= fold_build2 (MULT_EXPR
, llutype
, t3
,
792 single_nonrect_count
);
793 t3
= fold_build2 (TRUNC_DIV_EXPR
, llutype
, t3
,
794 build_int_cst (llutype
, 2));
795 t2
= fold_build2 (MULT_EXPR
, llutype
, t2
, t3
);
796 t
= fold_build2 (PLUS_EXPR
, llutype
, t
, t2
);
798 if (i
== single_nonrect
)
800 if (integer_zerop (t
) || TREE_CODE (t
) != INTEGER_CST
)
804 single_nonrect_count
= t
;
805 single_nonrect_cond_code
= loop
->cond_code
;
806 if (count
== NULL_TREE
)
807 count
= build_one_cst (llutype
);
810 else if (count
!= NULL_TREE
)
811 count
= fold_build2 (MULT_EXPR
, llutype
, count
, t
);
814 if (TREE_CODE (count
) != INTEGER_CST
)
817 else if (count
&& !integer_zerop (count
))
824 && (fd
->sched_kind
!= OMP_CLAUSE_SCHEDULE_STATIC
825 || fd
->have_ordered
))
827 if (!tree_int_cst_lt (count
, TYPE_MAX_VALUE (long_integer_type_node
)))
828 iter_type
= long_long_unsigned_type_node
;
830 iter_type
= long_integer_type_node
;
832 else if (collapse_iter
&& *collapse_iter
!= NULL
)
833 iter_type
= TREE_TYPE (*collapse_iter
);
834 fd
->iter_type
= iter_type
;
835 if (collapse_iter
&& *collapse_iter
== NULL
)
836 *collapse_iter
= create_tmp_var (iter_type
, ".iter");
837 if (collapse_count
&& *collapse_count
== NULL
)
841 *collapse_count
= fold_convert_loc (loc
, iter_type
, count
);
842 if (fd
->first_inner_iterations
&& fd
->factor
)
844 t
= make_tree_vec (4);
845 TREE_VEC_ELT (t
, 0) = *collapse_count
;
846 TREE_VEC_ELT (t
, 1) = fd
->first_inner_iterations
;
847 TREE_VEC_ELT (t
, 2) = fd
->factor
;
848 TREE_VEC_ELT (t
, 3) = fd
->adjn1
;
853 *collapse_count
= create_tmp_var (iter_type
, ".count");
856 if (fd
->collapse
> 1 || fd
->tiling
|| (fd
->ordered
&& loops
))
858 fd
->loop
.v
= *collapse_iter
;
859 fd
->loop
.n1
= build_int_cst (TREE_TYPE (fd
->loop
.v
), 0);
860 fd
->loop
.n2
= *collapse_count
;
861 if (TREE_CODE (fd
->loop
.n2
) == TREE_VEC
)
863 gcc_assert (fd
->non_rect
);
864 fd
->first_inner_iterations
= TREE_VEC_ELT (fd
->loop
.n2
, 1);
865 fd
->factor
= TREE_VEC_ELT (fd
->loop
.n2
, 2);
866 fd
->adjn1
= TREE_VEC_ELT (fd
->loop
.n2
, 3);
867 fd
->loop
.n2
= TREE_VEC_ELT (fd
->loop
.n2
, 0);
869 fd
->loop
.step
= build_int_cst (TREE_TYPE (fd
->loop
.v
), 1);
870 fd
->loop
.m1
= NULL_TREE
;
871 fd
->loop
.m2
= NULL_TREE
;
873 fd
->loop
.cond_code
= LT_EXPR
;
879 /* Build a call to GOMP_barrier. */
882 omp_build_barrier (tree lhs
)
884 tree fndecl
= builtin_decl_explicit (lhs
? BUILT_IN_GOMP_BARRIER_CANCEL
885 : BUILT_IN_GOMP_BARRIER
);
886 gcall
*g
= gimple_build_call (fndecl
, 0);
888 gimple_call_set_lhs (g
, lhs
);
892 /* Find OMP_FOR resp. OMP_SIMD with non-NULL OMP_FOR_INIT. Also, fill in pdata
893 array, pdata[0] non-NULL if there is anything non-trivial in between,
894 pdata[1] is address of OMP_PARALLEL in between if any, pdata[2] is address
895 of OMP_FOR in between if any and pdata[3] is address of the inner
899 find_combined_omp_for (tree
*tp
, int *walk_subtrees
, void *data
)
901 tree
**pdata
= (tree
**) data
;
903 switch (TREE_CODE (*tp
))
906 if (OMP_FOR_INIT (*tp
) != NULL_TREE
)
915 if (OMP_FOR_INIT (*tp
) != NULL_TREE
)
922 if (BIND_EXPR_VARS (*tp
)
923 || (BIND_EXPR_BLOCK (*tp
)
924 && BLOCK_VARS (BIND_EXPR_BLOCK (*tp
))))
929 if (!tsi_one_before_end_p (tsi_start (*tp
)))
933 case TRY_FINALLY_EXPR
:
947 /* Return maximum possible vectorization factor for the target. */
954 || !flag_tree_loop_optimize
955 || (!flag_tree_loop_vectorize
956 && OPTION_SET_P (flag_tree_loop_vectorize
)))
959 auto_vector_modes modes
;
960 targetm
.vectorize
.autovectorize_vector_modes (&modes
, true);
961 if (!modes
.is_empty ())
964 for (unsigned int i
= 0; i
< modes
.length (); ++i
)
965 /* The returned modes use the smallest element size (and thus
966 the largest nunits) for the vectorization approach that they
968 vf
= ordered_max (vf
, GET_MODE_NUNITS (modes
[i
]));
972 machine_mode vqimode
= targetm
.vectorize
.preferred_simd_mode (QImode
);
973 if (GET_MODE_CLASS (vqimode
) == MODE_VECTOR_INT
)
974 return GET_MODE_NUNITS (vqimode
);
979 /* Return maximum SIMT width if offloading may target SIMT hardware. */
982 omp_max_simt_vf (void)
986 if (ENABLE_OFFLOADING
)
987 for (const char *c
= getenv ("OFFLOAD_TARGET_NAMES"); c
;)
989 if (startswith (c
, "nvptx"))
991 else if ((c
= strchr (c
, ':')))
997 /* Store the construct selectors as tree codes from last to first,
998 return their number. */
1001 omp_constructor_traits_to_codes (tree ctx
, enum tree_code
*constructs
)
1003 int nconstructs
= list_length (ctx
);
1004 int i
= nconstructs
- 1;
1005 for (tree t2
= ctx
; t2
; t2
= TREE_CHAIN (t2
), i
--)
1007 const char *sel
= IDENTIFIER_POINTER (TREE_PURPOSE (t2
));
1008 if (!strcmp (sel
, "target"))
1009 constructs
[i
] = OMP_TARGET
;
1010 else if (!strcmp (sel
, "teams"))
1011 constructs
[i
] = OMP_TEAMS
;
1012 else if (!strcmp (sel
, "parallel"))
1013 constructs
[i
] = OMP_PARALLEL
;
1014 else if (!strcmp (sel
, "for") || !strcmp (sel
, "do"))
1015 constructs
[i
] = OMP_FOR
;
1016 else if (!strcmp (sel
, "simd"))
1017 constructs
[i
] = OMP_SIMD
;
1021 gcc_assert (i
== -1);
1025 /* Return true if PROP is possibly present in one of the offloading target's
1026 OpenMP contexts. The format of PROPS string is always offloading target's
1027 name terminated by '\0', followed by properties for that offloading
1028 target separated by '\0' and terminated by another '\0'. The strings
1029 are created from omp-device-properties installed files of all configured
1030 offloading targets. */
1033 omp_offload_device_kind_arch_isa (const char *props
, const char *prop
)
1035 const char *names
= getenv ("OFFLOAD_TARGET_NAMES");
1036 if (names
== NULL
|| *names
== '\0')
1038 while (*props
!= '\0')
1040 size_t name_len
= strlen (props
);
1041 bool matches
= false;
1042 for (const char *c
= names
; c
; )
1044 if (strncmp (props
, c
, name_len
) == 0
1045 && (c
[name_len
] == '\0'
1046 || c
[name_len
] == ':'
1047 || c
[name_len
] == '='))
1052 else if ((c
= strchr (c
, ':')))
1055 props
= props
+ name_len
+ 1;
1056 while (*props
!= '\0')
1058 if (matches
&& strcmp (props
, prop
) == 0)
1060 props
= strchr (props
, '\0') + 1;
1067 /* Return true if the current code location is or might be offloaded.
1068 Return true in declare target functions, or when nested in a target
1069 region or when unsure, return false otherwise. */
1072 omp_maybe_offloaded (void)
1074 if (!ENABLE_OFFLOADING
)
1076 const char *names
= getenv ("OFFLOAD_TARGET_NAMES");
1077 if (names
== NULL
|| *names
== '\0')
1080 if (symtab
->state
== PARSING
)
1083 if (cfun
&& cfun
->after_inlining
)
1085 if (current_function_decl
1086 && lookup_attribute ("omp declare target",
1087 DECL_ATTRIBUTES (current_function_decl
)))
1089 if (cfun
&& (cfun
->curr_properties
& PROP_gimple_any
) == 0)
1091 enum tree_code construct
= OMP_TARGET
;
1092 if (omp_construct_selector_matches (&construct
, 1, NULL
))
1099 /* Diagnose errors in an OpenMP context selector, return CTX if
1100 it is correct or error_mark_node otherwise. */
1103 omp_check_context_selector (location_t loc
, tree ctx
)
1105 /* Each trait-set-selector-name can only be specified once.
1106 There are just 4 set names. */
1107 for (tree t1
= ctx
; t1
; t1
= TREE_CHAIN (t1
))
1108 for (tree t2
= TREE_CHAIN (t1
); t2
; t2
= TREE_CHAIN (t2
))
1109 if (TREE_PURPOSE (t1
) == TREE_PURPOSE (t2
))
1111 error_at (loc
, "selector set %qs specified more than once",
1112 IDENTIFIER_POINTER (TREE_PURPOSE (t1
)));
1113 return error_mark_node
;
1115 for (tree t
= ctx
; t
; t
= TREE_CHAIN (t
))
1117 /* Each trait-selector-name can only be specified once. */
1118 if (list_length (TREE_VALUE (t
)) < 5)
1120 for (tree t1
= TREE_VALUE (t
); t1
; t1
= TREE_CHAIN (t1
))
1121 for (tree t2
= TREE_CHAIN (t1
); t2
; t2
= TREE_CHAIN (t2
))
1122 if (TREE_PURPOSE (t1
) == TREE_PURPOSE (t2
))
1125 "selector %qs specified more than once in set %qs",
1126 IDENTIFIER_POINTER (TREE_PURPOSE (t1
)),
1127 IDENTIFIER_POINTER (TREE_PURPOSE (t
)));
1128 return error_mark_node
;
1133 hash_set
<tree
> pset
;
1134 for (tree t1
= TREE_VALUE (t
); t1
; t1
= TREE_CHAIN (t1
))
1135 if (pset
.add (TREE_PURPOSE (t1
)))
1138 "selector %qs specified more than once in set %qs",
1139 IDENTIFIER_POINTER (TREE_PURPOSE (t1
)),
1140 IDENTIFIER_POINTER (TREE_PURPOSE (t
)));
1141 return error_mark_node
;
1145 static const char *const kind
[] = {
1146 "host", "nohost", "cpu", "gpu", "fpga", "any", NULL
};
1147 static const char *const vendor
[] = {
1148 "amd", "arm", "bsc", "cray", "fujitsu", "gnu", "ibm", "intel",
1149 "llvm", "nvidia", "pgi", "ti", "unknown", NULL
};
1150 static const char *const extension
[] = { NULL
};
1151 static const char *const atomic_default_mem_order
[] = {
1152 "seq_cst", "relaxed", "acq_rel", NULL
};
1153 struct known_properties
{ const char *set
; const char *selector
;
1154 const char *const *props
; };
1155 known_properties props
[] = {
1156 { "device", "kind", kind
},
1157 { "implementation", "vendor", vendor
},
1158 { "implementation", "extension", extension
},
1159 { "implementation", "atomic_default_mem_order",
1160 atomic_default_mem_order
} };
1161 for (tree t1
= TREE_VALUE (t
); t1
; t1
= TREE_CHAIN (t1
))
1162 for (unsigned i
= 0; i
< ARRAY_SIZE (props
); i
++)
1163 if (!strcmp (IDENTIFIER_POINTER (TREE_PURPOSE (t1
)),
1165 && !strcmp (IDENTIFIER_POINTER (TREE_PURPOSE (t
)),
1167 for (tree t2
= TREE_VALUE (t1
); t2
; t2
= TREE_CHAIN (t2
))
1168 for (unsigned j
= 0; ; j
++)
1170 if (props
[i
].props
[j
] == NULL
)
1172 if (TREE_PURPOSE (t2
)
1173 && !strcmp (IDENTIFIER_POINTER (TREE_PURPOSE (t2
)),
1176 if (props
[i
].props
== atomic_default_mem_order
)
1179 "incorrect property %qs of %qs selector",
1180 IDENTIFIER_POINTER (TREE_PURPOSE (t2
)),
1181 "atomic_default_mem_order");
1182 return error_mark_node
;
1184 else if (TREE_PURPOSE (t2
))
1186 "unknown property %qs of %qs selector",
1187 IDENTIFIER_POINTER (TREE_PURPOSE (t2
)),
1191 "unknown property %qE of %qs selector",
1192 TREE_VALUE (t2
), props
[i
].selector
);
1195 else if (TREE_PURPOSE (t2
) == NULL_TREE
)
1197 const char *str
= TREE_STRING_POINTER (TREE_VALUE (t2
));
1198 if (!strcmp (str
, props
[i
].props
[j
])
1199 && ((size_t) TREE_STRING_LENGTH (TREE_VALUE (t2
))
1200 == strlen (str
) + (lang_GNU_Fortran () ? 0 : 1)))
1203 else if (!strcmp (IDENTIFIER_POINTER (TREE_PURPOSE (t2
)),
1212 /* Register VARIANT as variant of some base function marked with
1213 #pragma omp declare variant. CONSTRUCT is corresponding construct
1217 omp_mark_declare_variant (location_t loc
, tree variant
, tree construct
)
1219 tree attr
= lookup_attribute ("omp declare variant variant",
1220 DECL_ATTRIBUTES (variant
));
1221 if (attr
== NULL_TREE
)
1223 attr
= tree_cons (get_identifier ("omp declare variant variant"),
1224 unshare_expr (construct
),
1225 DECL_ATTRIBUTES (variant
));
1226 DECL_ATTRIBUTES (variant
) = attr
;
1229 if ((TREE_VALUE (attr
) != NULL_TREE
) != (construct
!= NULL_TREE
)
1230 || (construct
!= NULL_TREE
1231 && omp_context_selector_set_compare ("construct", TREE_VALUE (attr
),
1233 error_at (loc
, "%qD used as a variant with incompatible %<construct%> "
1234 "selector sets", variant
);
1238 /* Return a name from PROP, a property in selectors accepting
1242 omp_context_name_list_prop (tree prop
)
1244 if (TREE_PURPOSE (prop
))
1245 return IDENTIFIER_POINTER (TREE_PURPOSE (prop
));
1248 const char *ret
= TREE_STRING_POINTER (TREE_VALUE (prop
));
1249 if ((size_t) TREE_STRING_LENGTH (TREE_VALUE (prop
))
1250 == strlen (ret
) + (lang_GNU_Fortran () ? 0 : 1))
1256 /* Return 1 if context selector matches the current OpenMP context, 0
1257 if it does not and -1 if it is unknown and need to be determined later.
1258 Some properties can be checked right away during parsing (this routine),
1259 others need to wait until the whole TU is parsed, others need to wait until
1260 IPA, others until vectorization. */
1263 omp_context_selector_matches (tree ctx
)
1266 for (tree t1
= ctx
; t1
; t1
= TREE_CHAIN (t1
))
1268 char set
= IDENTIFIER_POINTER (TREE_PURPOSE (t1
))[0];
1271 /* For now, ignore the construct set. While something can be
1272 determined already during parsing, we don't know until end of TU
1273 whether additional constructs aren't added through declare variant
1274 unless "omp declare variant variant" attribute exists already
1275 (so in most of the cases), and we'd need to maintain set of
1276 surrounding OpenMP constructs, which is better handled during
1278 if (symtab
->state
== PARSING
)
1284 enum tree_code constructs
[5];
1286 = omp_constructor_traits_to_codes (TREE_VALUE (t1
), constructs
);
1288 if (cfun
&& (cfun
->curr_properties
& PROP_gimple_any
) != 0)
1290 if (!cfun
->after_inlining
)
1296 for (i
= 0; i
< nconstructs
; ++i
)
1297 if (constructs
[i
] == OMP_SIMD
)
1299 if (i
< nconstructs
)
1304 /* If there is no simd, assume it is ok after IPA,
1305 constructs should have been checked before. */
1309 int r
= omp_construct_selector_matches (constructs
, nconstructs
,
1317 for (tree t2
= TREE_VALUE (t1
); t2
; t2
= TREE_CHAIN (t2
))
1319 const char *sel
= IDENTIFIER_POINTER (TREE_PURPOSE (t2
));
1323 if (set
== 'i' && !strcmp (sel
, "vendor"))
1324 for (tree t3
= TREE_VALUE (t2
); t3
; t3
= TREE_CHAIN (t3
))
1326 const char *prop
= omp_context_name_list_prop (t3
);
1329 if ((!strcmp (prop
, " score") && TREE_PURPOSE (t3
))
1330 || !strcmp (prop
, "gnu"))
1336 if (set
== 'i' && !strcmp (sel
, "extension"))
1337 /* We don't support any extensions right now. */
1341 if (set
== 'i' && !strcmp (sel
, "atomic_default_mem_order"))
1343 if (cfun
&& (cfun
->curr_properties
& PROP_gimple_any
) != 0)
1346 enum omp_memory_order omo
1347 = ((enum omp_memory_order
)
1349 & OMP_REQUIRES_ATOMIC_DEFAULT_MEM_ORDER
));
1350 if (omo
== OMP_MEMORY_ORDER_UNSPECIFIED
)
1352 /* We don't know yet, until end of TU. */
1353 if (symtab
->state
== PARSING
)
1359 omo
= OMP_MEMORY_ORDER_RELAXED
;
1361 tree t3
= TREE_VALUE (t2
);
1362 const char *prop
= IDENTIFIER_POINTER (TREE_PURPOSE (t3
));
1363 if (!strcmp (prop
, " score"))
1365 t3
= TREE_CHAIN (t3
);
1366 prop
= IDENTIFIER_POINTER (TREE_PURPOSE (t3
));
1368 if (!strcmp (prop
, "relaxed")
1369 && omo
!= OMP_MEMORY_ORDER_RELAXED
)
1371 else if (!strcmp (prop
, "seq_cst")
1372 && omo
!= OMP_MEMORY_ORDER_SEQ_CST
)
1374 else if (!strcmp (prop
, "acq_rel")
1375 && omo
!= OMP_MEMORY_ORDER_ACQ_REL
)
1378 if (set
== 'd' && !strcmp (sel
, "arch"))
1379 for (tree t3
= TREE_VALUE (t2
); t3
; t3
= TREE_CHAIN (t3
))
1381 const char *arch
= omp_context_name_list_prop (t3
);
1385 if (targetm
.omp
.device_kind_arch_isa
!= NULL
)
1386 r
= targetm
.omp
.device_kind_arch_isa (omp_device_arch
,
1388 if (r
== 0 || (r
== -1 && symtab
->state
!= PARSING
))
1390 /* If we are or might be in a target region or
1391 declare target function, need to take into account
1392 also offloading values. */
1393 if (!omp_maybe_offloaded ())
1395 if (ENABLE_OFFLOADING
)
1397 const char *arches
= omp_offload_device_arch
;
1398 if (omp_offload_device_kind_arch_isa (arches
,
1409 /* If arch matches on the host, it still might not match
1410 in the offloading region. */
1411 else if (omp_maybe_offloaded ())
1416 if (set
== 'i' && !strcmp (sel
, "unified_address"))
1418 if (cfun
&& (cfun
->curr_properties
& PROP_gimple_any
) != 0)
1421 if ((omp_requires_mask
& OMP_REQUIRES_UNIFIED_ADDRESS
) == 0)
1423 if (symtab
->state
== PARSING
)
1430 if (set
== 'i' && !strcmp (sel
, "unified_shared_memory"))
1432 if (cfun
&& (cfun
->curr_properties
& PROP_gimple_any
) != 0)
1435 if ((omp_requires_mask
1436 & OMP_REQUIRES_UNIFIED_SHARED_MEMORY
) == 0)
1438 if (symtab
->state
== PARSING
)
1447 if (set
== 'i' && !strcmp (sel
, "dynamic_allocators"))
1449 if (cfun
&& (cfun
->curr_properties
& PROP_gimple_any
) != 0)
1452 if ((omp_requires_mask
1453 & OMP_REQUIRES_DYNAMIC_ALLOCATORS
) == 0)
1455 if (symtab
->state
== PARSING
)
1464 if (set
== 'i' && !strcmp (sel
, "reverse_offload"))
1466 if (cfun
&& (cfun
->curr_properties
& PROP_gimple_any
) != 0)
1469 if ((omp_requires_mask
& OMP_REQUIRES_REVERSE_OFFLOAD
) == 0)
1471 if (symtab
->state
== PARSING
)
1480 if (set
== 'd' && !strcmp (sel
, "kind"))
1481 for (tree t3
= TREE_VALUE (t2
); t3
; t3
= TREE_CHAIN (t3
))
1483 const char *prop
= omp_context_name_list_prop (t3
);
1486 if (!strcmp (prop
, "any"))
1488 if (!strcmp (prop
, "host"))
1490 #ifdef ACCEL_COMPILER
1493 if (omp_maybe_offloaded ())
1498 if (!strcmp (prop
, "nohost"))
1500 #ifndef ACCEL_COMPILER
1501 if (omp_maybe_offloaded ())
1509 if (targetm
.omp
.device_kind_arch_isa
!= NULL
)
1510 r
= targetm
.omp
.device_kind_arch_isa (omp_device_kind
,
1513 r
= strcmp (prop
, "cpu") == 0;
1514 if (r
== 0 || (r
== -1 && symtab
->state
!= PARSING
))
1516 /* If we are or might be in a target region or
1517 declare target function, need to take into account
1518 also offloading values. */
1519 if (!omp_maybe_offloaded ())
1521 if (ENABLE_OFFLOADING
)
1523 const char *kinds
= omp_offload_device_kind
;
1524 if (omp_offload_device_kind_arch_isa (kinds
, prop
))
1534 /* If kind matches on the host, it still might not match
1535 in the offloading region. */
1536 else if (omp_maybe_offloaded ())
1541 if (set
== 'd' && !strcmp (sel
, "isa"))
1542 for (tree t3
= TREE_VALUE (t2
); t3
; t3
= TREE_CHAIN (t3
))
1544 const char *isa
= omp_context_name_list_prop (t3
);
1548 if (targetm
.omp
.device_kind_arch_isa
!= NULL
)
1549 r
= targetm
.omp
.device_kind_arch_isa (omp_device_isa
,
1551 if (r
== 0 || (r
== -1 && symtab
->state
!= PARSING
))
1553 /* If isa is valid on the target, but not in the
1554 current function and current function has
1555 #pragma omp declare simd on it, some simd clones
1556 might have the isa added later on. */
1558 && targetm
.simd_clone
.compute_vecsize_and_simdlen
1559 && (cfun
== NULL
|| !cfun
->after_inlining
))
1562 = DECL_ATTRIBUTES (current_function_decl
);
1563 if (lookup_attribute ("omp declare simd", attrs
))
1569 /* If we are or might be in a target region or
1570 declare target function, need to take into account
1571 also offloading values. */
1572 if (!omp_maybe_offloaded ())
1574 if (ENABLE_OFFLOADING
)
1576 const char *isas
= omp_offload_device_isa
;
1577 if (omp_offload_device_kind_arch_isa (isas
, isa
))
1587 /* If isa matches on the host, it still might not match
1588 in the offloading region. */
1589 else if (omp_maybe_offloaded ())
1594 if (set
== 'u' && !strcmp (sel
, "condition"))
1595 for (tree t3
= TREE_VALUE (t2
); t3
; t3
= TREE_CHAIN (t3
))
1596 if (TREE_PURPOSE (t3
) == NULL_TREE
)
1598 if (integer_zerop (TREE_VALUE (t3
)))
1600 if (integer_nonzerop (TREE_VALUE (t3
)))
1613 /* Compare construct={simd} CLAUSES1 with CLAUSES2, return 0/-1/1/2 as
1614 in omp_context_selector_set_compare. */
1617 omp_construct_simd_compare (tree clauses1
, tree clauses2
)
1619 if (clauses1
== NULL_TREE
)
1620 return clauses2
== NULL_TREE
? 0 : -1;
1621 if (clauses2
== NULL_TREE
)
1625 struct declare_variant_simd_data
{
1626 bool inbranch
, notinbranch
;
1628 auto_vec
<tree
,16> data_sharing
;
1629 auto_vec
<tree
,16> aligned
;
1630 declare_variant_simd_data ()
1631 : inbranch(false), notinbranch(false), simdlen(NULL_TREE
) {}
1634 for (i
= 0; i
< 2; i
++)
1635 for (tree c
= i
? clauses2
: clauses1
; c
; c
= OMP_CLAUSE_CHAIN (c
))
1638 switch (OMP_CLAUSE_CODE (c
))
1640 case OMP_CLAUSE_INBRANCH
:
1641 data
[i
].inbranch
= true;
1643 case OMP_CLAUSE_NOTINBRANCH
:
1644 data
[i
].notinbranch
= true;
1646 case OMP_CLAUSE_SIMDLEN
:
1647 data
[i
].simdlen
= OMP_CLAUSE_SIMDLEN_EXPR (c
);
1649 case OMP_CLAUSE_UNIFORM
:
1650 case OMP_CLAUSE_LINEAR
:
1651 v
= &data
[i
].data_sharing
;
1653 case OMP_CLAUSE_ALIGNED
:
1654 v
= &data
[i
].aligned
;
1659 unsigned HOST_WIDE_INT argno
= tree_to_uhwi (OMP_CLAUSE_DECL (c
));
1660 if (argno
>= v
->length ())
1661 v
->safe_grow_cleared (argno
+ 1, true);
1664 /* Here, r is used as a bitmask, 2 is set if CLAUSES1 has something
1665 CLAUSES2 doesn't, 1 is set if CLAUSES2 has something CLAUSES1
1666 doesn't. Thus, r == 3 implies return value 2, r == 1 implies
1667 -1, r == 2 implies 1 and r == 0 implies 0. */
1668 if (data
[0].inbranch
!= data
[1].inbranch
)
1669 r
|= data
[0].inbranch
? 2 : 1;
1670 if (data
[0].notinbranch
!= data
[1].notinbranch
)
1671 r
|= data
[0].notinbranch
? 2 : 1;
1672 if (!simple_cst_equal (data
[0].simdlen
, data
[1].simdlen
))
1674 if (data
[0].simdlen
&& data
[1].simdlen
)
1676 r
|= data
[0].simdlen
? 2 : 1;
1678 if (data
[0].data_sharing
.length () < data
[1].data_sharing
.length ()
1679 || data
[0].aligned
.length () < data
[1].aligned
.length ())
1682 FOR_EACH_VEC_ELT (data
[0].data_sharing
, i
, c1
)
1684 c2
= (i
< data
[1].data_sharing
.length ()
1685 ? data
[1].data_sharing
[i
] : NULL_TREE
);
1686 if ((c1
== NULL_TREE
) != (c2
== NULL_TREE
))
1688 r
|= c1
!= NULL_TREE
? 2 : 1;
1691 if (c1
== NULL_TREE
)
1693 if (OMP_CLAUSE_CODE (c1
) != OMP_CLAUSE_CODE (c2
))
1695 if (OMP_CLAUSE_CODE (c1
) != OMP_CLAUSE_LINEAR
)
1697 if (OMP_CLAUSE_LINEAR_VARIABLE_STRIDE (c1
)
1698 != OMP_CLAUSE_LINEAR_VARIABLE_STRIDE (c2
))
1700 if (OMP_CLAUSE_LINEAR_KIND (c1
) != OMP_CLAUSE_LINEAR_KIND (c2
))
1702 if (!simple_cst_equal (OMP_CLAUSE_LINEAR_STEP (c1
),
1703 OMP_CLAUSE_LINEAR_STEP (c2
)))
1706 FOR_EACH_VEC_ELT (data
[0].aligned
, i
, c1
)
1708 c2
= i
< data
[1].aligned
.length () ? data
[1].aligned
[i
] : NULL_TREE
;
1709 if ((c1
== NULL_TREE
) != (c2
== NULL_TREE
))
1711 r
|= c1
!= NULL_TREE
? 2 : 1;
1714 if (c1
== NULL_TREE
)
1716 if (!simple_cst_equal (OMP_CLAUSE_ALIGNED_ALIGNMENT (c1
),
1717 OMP_CLAUSE_ALIGNED_ALIGNMENT (c2
)))
1726 default: gcc_unreachable ();
1730 /* Compare properties of selectors SEL from SET other than construct.
1731 Return 0/-1/1/2 as in omp_context_selector_set_compare.
1732 Unlike set names or selector names, properties can have duplicates. */
1735 omp_context_selector_props_compare (const char *set
, const char *sel
,
1736 tree ctx1
, tree ctx2
)
1739 for (int pass
= 0; pass
< 2; pass
++)
1740 for (tree t1
= pass
? ctx2
: ctx1
; t1
; t1
= TREE_CHAIN (t1
))
1743 for (t2
= pass
? ctx1
: ctx2
; t2
; t2
= TREE_CHAIN (t2
))
1744 if (TREE_PURPOSE (t1
) == TREE_PURPOSE (t2
))
1746 if (TREE_PURPOSE (t1
) == NULL_TREE
)
1748 if (set
[0] == 'u' && strcmp (sel
, "condition") == 0)
1750 if (integer_zerop (TREE_VALUE (t1
))
1751 != integer_zerop (TREE_VALUE (t2
)))
1755 if (simple_cst_equal (TREE_VALUE (t1
), TREE_VALUE (t2
)))
1758 else if (strcmp (IDENTIFIER_POINTER (TREE_PURPOSE (t1
)),
1761 if (!simple_cst_equal (TREE_VALUE (t1
), TREE_VALUE (t2
)))
1768 else if (TREE_PURPOSE (t1
)
1769 && TREE_PURPOSE (t2
) == NULL_TREE
1770 && TREE_CODE (TREE_VALUE (t2
)) == STRING_CST
)
1772 const char *p1
= omp_context_name_list_prop (t1
);
1773 const char *p2
= omp_context_name_list_prop (t2
);
1775 && strcmp (p1
, p2
) == 0
1776 && strcmp (p1
, " score"))
1779 else if (TREE_PURPOSE (t1
) == NULL_TREE
1780 && TREE_PURPOSE (t2
)
1781 && TREE_CODE (TREE_VALUE (t1
)) == STRING_CST
)
1783 const char *p1
= omp_context_name_list_prop (t1
);
1784 const char *p2
= omp_context_name_list_prop (t2
);
1786 && strcmp (p1
, p2
) == 0
1787 && strcmp (p1
, " score"))
1790 if (t2
== NULL_TREE
)
1792 int r
= pass
? -1 : 1;
1793 if (ret
&& ret
!= r
)
1807 /* Compare single context selector sets CTX1 and CTX2 with SET name.
1808 Return 0 if CTX1 is equal to CTX2,
1809 -1 if CTX1 is a strict subset of CTX2,
1810 1 if CTX2 is a strict subset of CTX1, or
1811 2 if neither context is a subset of another one. */
1814 omp_context_selector_set_compare (const char *set
, tree ctx1
, tree ctx2
)
1816 bool swapped
= false;
1818 int len1
= list_length (ctx1
);
1819 int len2
= list_length (ctx2
);
1824 std::swap (ctx1
, ctx2
);
1825 std::swap (len1
, len2
);
1831 tree simd
= get_identifier ("simd");
1832 /* Handle construct set specially. In this case the order
1833 of the selector matters too. */
1834 for (t1
= ctx1
; t1
; t1
= TREE_CHAIN (t1
))
1835 if (TREE_PURPOSE (t1
) == TREE_PURPOSE (t2
))
1838 if (TREE_PURPOSE (t1
) == simd
)
1839 r
= omp_construct_simd_compare (TREE_VALUE (t1
),
1841 if (r
== 2 || (ret
&& r
&& (ret
< 0) != (r
< 0)))
1845 t2
= TREE_CHAIN (t2
);
1846 if (t2
== NULL_TREE
)
1848 t1
= TREE_CHAIN (t1
);
1856 if (t2
!= NULL_TREE
)
1858 if (t1
!= NULL_TREE
)
1866 return swapped
? -ret
: ret
;
1868 for (tree t1
= ctx1
; t1
; t1
= TREE_CHAIN (t1
))
1871 for (t2
= ctx2
; t2
; t2
= TREE_CHAIN (t2
))
1872 if (TREE_PURPOSE (t1
) == TREE_PURPOSE (t2
))
1874 const char *sel
= IDENTIFIER_POINTER (TREE_PURPOSE (t1
));
1875 int r
= omp_context_selector_props_compare (set
, sel
,
1878 if (r
== 2 || (ret
&& r
&& (ret
< 0) != (r
< 0)))
1885 if (t2
== NULL_TREE
)
1896 return swapped
? -ret
: ret
;
1899 /* Compare whole context selector specification CTX1 and CTX2.
1900 Return 0 if CTX1 is equal to CTX2,
1901 -1 if CTX1 is a strict subset of CTX2,
1902 1 if CTX2 is a strict subset of CTX1, or
1903 2 if neither context is a subset of another one. */
1906 omp_context_selector_compare (tree ctx1
, tree ctx2
)
1908 bool swapped
= false;
1910 int len1
= list_length (ctx1
);
1911 int len2
= list_length (ctx2
);
1916 std::swap (ctx1
, ctx2
);
1917 std::swap (len1
, len2
);
1919 for (tree t1
= ctx1
; t1
; t1
= TREE_CHAIN (t1
))
1922 for (t2
= ctx2
; t2
; t2
= TREE_CHAIN (t2
))
1923 if (TREE_PURPOSE (t1
) == TREE_PURPOSE (t2
))
1925 const char *set
= IDENTIFIER_POINTER (TREE_PURPOSE (t1
));
1926 int r
= omp_context_selector_set_compare (set
, TREE_VALUE (t1
),
1928 if (r
== 2 || (ret
&& r
&& (ret
< 0) != (r
< 0)))
1935 if (t2
== NULL_TREE
)
1946 return swapped
? -ret
: ret
;
1949 /* From context selector CTX, return trait-selector with name SEL in
1950 trait-selector-set with name SET if any, or NULL_TREE if not found.
1951 If SEL is NULL, return the list of trait-selectors in SET. */
1954 omp_get_context_selector (tree ctx
, const char *set
, const char *sel
)
1956 tree setid
= get_identifier (set
);
1957 tree selid
= sel
? get_identifier (sel
) : NULL_TREE
;
1958 for (tree t1
= ctx
; t1
; t1
= TREE_CHAIN (t1
))
1959 if (TREE_PURPOSE (t1
) == setid
)
1962 return TREE_VALUE (t1
);
1963 for (tree t2
= TREE_VALUE (t1
); t2
; t2
= TREE_CHAIN (t2
))
1964 if (TREE_PURPOSE (t2
) == selid
)
1970 /* Compute *SCORE for context selector CTX. Return true if the score
1971 would be different depending on whether it is a declare simd clone or
1972 not. DECLARE_SIMD should be true for the case when it would be
1973 a declare simd clone. */
1976 omp_context_compute_score (tree ctx
, widest_int
*score
, bool declare_simd
)
1978 tree construct
= omp_get_context_selector (ctx
, "construct", NULL
);
1979 bool has_kind
= omp_get_context_selector (ctx
, "device", "kind");
1980 bool has_arch
= omp_get_context_selector (ctx
, "device", "arch");
1981 bool has_isa
= omp_get_context_selector (ctx
, "device", "isa");
1984 for (tree t1
= ctx
; t1
; t1
= TREE_CHAIN (t1
))
1985 if (TREE_VALUE (t1
) != construct
)
1986 for (tree t2
= TREE_VALUE (t1
); t2
; t2
= TREE_CHAIN (t2
))
1987 if (tree t3
= TREE_VALUE (t2
))
1988 if (TREE_PURPOSE (t3
)
1989 && strcmp (IDENTIFIER_POINTER (TREE_PURPOSE (t3
)), " score") == 0
1990 && TREE_CODE (TREE_VALUE (t3
)) == INTEGER_CST
)
1991 *score
+= wi::to_widest (TREE_VALUE (t3
));
1992 if (construct
|| has_kind
|| has_arch
|| has_isa
)
1995 enum tree_code constructs
[5];
1996 int nconstructs
= 0;
1998 nconstructs
= omp_constructor_traits_to_codes (construct
, constructs
);
1999 if (omp_construct_selector_matches (constructs
, nconstructs
, scores
)
2002 int b
= declare_simd
? nconstructs
+ 1 : 0;
2003 if (scores
[b
+ nconstructs
] + 4U < score
->get_precision ())
2005 for (int n
= 0; n
< nconstructs
; ++n
)
2007 if (scores
[b
+ n
] < 0)
2012 *score
+= wi::shifted_mask
<widest_int
> (scores
[b
+ n
], 1, false);
2015 *score
+= wi::shifted_mask
<widest_int
> (scores
[b
+ nconstructs
],
2018 *score
+= wi::shifted_mask
<widest_int
> (scores
[b
+ nconstructs
] + 1,
2021 *score
+= wi::shifted_mask
<widest_int
> (scores
[b
+ nconstructs
] + 2,
2024 else /* FIXME: Implement this. */
2030 /* Class describing a single variant. */
2031 struct GTY(()) omp_declare_variant_entry
{
2032 /* NODE of the variant. */
2033 cgraph_node
*variant
;
2034 /* Score if not in declare simd clone. */
2036 /* Score if in declare simd clone. */
2037 widest_int score_in_declare_simd_clone
;
2038 /* Context selector for the variant. */
2040 /* True if the context selector is known to match already. */
2044 /* Class describing a function with variants. */
2045 struct GTY((for_user
)) omp_declare_variant_base_entry
{
2046 /* NODE of the base function. */
2048 /* NODE of the artificial function created for the deferred variant
2051 /* Vector of the variants. */
2052 vec
<omp_declare_variant_entry
, va_gc
> *variants
;
2055 struct omp_declare_variant_hasher
2056 : ggc_ptr_hash
<omp_declare_variant_base_entry
> {
2057 static hashval_t
hash (omp_declare_variant_base_entry
*);
2058 static bool equal (omp_declare_variant_base_entry
*,
2059 omp_declare_variant_base_entry
*);
2063 omp_declare_variant_hasher::hash (omp_declare_variant_base_entry
*x
)
2065 inchash::hash hstate
;
2066 hstate
.add_int (DECL_UID (x
->base
->decl
));
2067 hstate
.add_int (x
->variants
->length ());
2068 omp_declare_variant_entry
*variant
;
2070 FOR_EACH_VEC_SAFE_ELT (x
->variants
, i
, variant
)
2072 hstate
.add_int (DECL_UID (variant
->variant
->decl
));
2073 hstate
.add_wide_int (variant
->score
);
2074 hstate
.add_wide_int (variant
->score_in_declare_simd_clone
);
2075 hstate
.add_ptr (variant
->ctx
);
2076 hstate
.add_int (variant
->matches
);
2078 return hstate
.end ();
2082 omp_declare_variant_hasher::equal (omp_declare_variant_base_entry
*x
,
2083 omp_declare_variant_base_entry
*y
)
2085 if (x
->base
!= y
->base
2086 || x
->variants
->length () != y
->variants
->length ())
2088 omp_declare_variant_entry
*variant
;
2090 FOR_EACH_VEC_SAFE_ELT (x
->variants
, i
, variant
)
2091 if (variant
->variant
!= (*y
->variants
)[i
].variant
2092 || variant
->score
!= (*y
->variants
)[i
].score
2093 || (variant
->score_in_declare_simd_clone
2094 != (*y
->variants
)[i
].score_in_declare_simd_clone
)
2095 || variant
->ctx
!= (*y
->variants
)[i
].ctx
2096 || variant
->matches
!= (*y
->variants
)[i
].matches
)
2101 static GTY(()) hash_table
<omp_declare_variant_hasher
> *omp_declare_variants
;
2103 struct omp_declare_variant_alt_hasher
2104 : ggc_ptr_hash
<omp_declare_variant_base_entry
> {
2105 static hashval_t
hash (omp_declare_variant_base_entry
*);
2106 static bool equal (omp_declare_variant_base_entry
*,
2107 omp_declare_variant_base_entry
*);
2111 omp_declare_variant_alt_hasher::hash (omp_declare_variant_base_entry
*x
)
2113 return DECL_UID (x
->node
->decl
);
2117 omp_declare_variant_alt_hasher::equal (omp_declare_variant_base_entry
*x
,
2118 omp_declare_variant_base_entry
*y
)
2120 return x
->node
== y
->node
;
2123 static GTY(()) hash_table
<omp_declare_variant_alt_hasher
>
2124 *omp_declare_variant_alt
;
2126 /* Try to resolve declare variant after gimplification. */
2129 omp_resolve_late_declare_variant (tree alt
)
2131 cgraph_node
*node
= cgraph_node::get (alt
);
2132 cgraph_node
*cur_node
= cgraph_node::get (cfun
->decl
);
2134 || !node
->declare_variant_alt
2135 || !cfun
->after_inlining
)
2138 omp_declare_variant_base_entry entry
;
2141 entry
.variants
= NULL
;
2142 omp_declare_variant_base_entry
*entryp
2143 = omp_declare_variant_alt
->find_with_hash (&entry
, DECL_UID (alt
));
2146 omp_declare_variant_entry
*varentry1
, *varentry2
;
2147 auto_vec
<bool, 16> matches
;
2148 unsigned int nmatches
= 0;
2149 FOR_EACH_VEC_SAFE_ELT (entryp
->variants
, i
, varentry1
)
2151 if (varentry1
->matches
)
2153 /* This has been checked to be ok already. */
2154 matches
.safe_push (true);
2158 switch (omp_context_selector_matches (varentry1
->ctx
))
2161 matches
.safe_push (false);
2166 matches
.safe_push (true);
2173 return entryp
->base
->decl
;
2175 /* A context selector that is a strict subset of another context selector
2176 has a score of zero. */
2177 FOR_EACH_VEC_SAFE_ELT (entryp
->variants
, i
, varentry1
)
2181 vec_safe_iterate (entryp
->variants
, j
, &varentry2
); ++j
)
2184 int r
= omp_context_selector_compare (varentry1
->ctx
,
2188 /* ctx1 is a strict subset of ctx2, ignore ctx1. */
2193 /* ctx2 is a strict subset of ctx1, remove ctx2. */
2198 widest_int max_score
= -1;
2200 FOR_EACH_VEC_SAFE_ELT (entryp
->variants
, i
, varentry1
)
2204 = (cur_node
->simdclone
? varentry1
->score_in_declare_simd_clone
2205 : varentry1
->score
);
2206 if (score
> max_score
)
2209 varentry2
= varentry1
;
2212 return varentry2
->variant
->decl
;
2215 /* Hook to adjust hash tables on cgraph_node removal. */
2218 omp_declare_variant_remove_hook (struct cgraph_node
*node
, void *)
2220 if (!node
->declare_variant_alt
)
2223 /* Drop this hash table completely. */
2224 omp_declare_variants
= NULL
;
2225 /* And remove node from the other hash table. */
2226 if (omp_declare_variant_alt
)
2228 omp_declare_variant_base_entry entry
;
2231 entry
.variants
= NULL
;
2232 omp_declare_variant_alt
->remove_elt_with_hash (&entry
,
2233 DECL_UID (node
->decl
));
2237 /* Try to resolve declare variant, return the variant decl if it should
2238 be used instead of base, or base otherwise. */
2241 omp_resolve_declare_variant (tree base
)
2243 tree variant1
= NULL_TREE
, variant2
= NULL_TREE
;
2244 if (cfun
&& (cfun
->curr_properties
& PROP_gimple_any
) != 0)
2245 return omp_resolve_late_declare_variant (base
);
2247 auto_vec
<tree
, 16> variants
;
2248 auto_vec
<bool, 16> defer
;
2249 bool any_deferred
= false;
2250 for (tree attr
= DECL_ATTRIBUTES (base
); attr
; attr
= TREE_CHAIN (attr
))
2252 attr
= lookup_attribute ("omp declare variant base", attr
);
2253 if (attr
== NULL_TREE
)
2255 if (TREE_CODE (TREE_PURPOSE (TREE_VALUE (attr
))) != FUNCTION_DECL
)
2257 cgraph_node
*node
= cgraph_node::get (base
);
2258 /* If this is already a magic decl created by this function,
2259 don't process it again. */
2260 if (node
&& node
->declare_variant_alt
)
2262 switch (omp_context_selector_matches (TREE_VALUE (TREE_VALUE (attr
))))
2265 /* No match, ignore. */
2268 /* Needs to be deferred. */
2269 any_deferred
= true;
2270 variants
.safe_push (attr
);
2271 defer
.safe_push (true);
2274 variants
.safe_push (attr
);
2275 defer
.safe_push (false);
2279 if (variants
.length () == 0)
2284 widest_int max_score1
= 0;
2285 widest_int max_score2
= 0;
2289 omp_declare_variant_base_entry entry
;
2290 entry
.base
= cgraph_node::get_create (base
);
2292 vec_alloc (entry
.variants
, variants
.length ());
2293 FOR_EACH_VEC_ELT (variants
, i
, attr1
)
2298 tree ctx
= TREE_VALUE (TREE_VALUE (attr1
));
2299 need_two
= omp_context_compute_score (ctx
, &score1
, false);
2301 omp_context_compute_score (ctx
, &score2
, true);
2307 max_score1
= score1
;
2308 max_score2
= score2
;
2317 if (max_score1
== score1
)
2318 variant1
= NULL_TREE
;
2319 else if (score1
> max_score1
)
2321 max_score1
= score1
;
2322 variant1
= defer
[i
] ? NULL_TREE
: attr1
;
2324 if (max_score2
== score2
)
2325 variant2
= NULL_TREE
;
2326 else if (score2
> max_score2
)
2328 max_score2
= score2
;
2329 variant2
= defer
[i
] ? NULL_TREE
: attr1
;
2332 omp_declare_variant_entry varentry
;
2334 = cgraph_node::get_create (TREE_PURPOSE (TREE_VALUE (attr1
)));
2335 varentry
.score
= score1
;
2336 varentry
.score_in_declare_simd_clone
= score2
;
2338 varentry
.matches
= !defer
[i
];
2339 entry
.variants
->quick_push (varentry
);
2342 /* If there is a clear winner variant with the score which is not
2343 deferred, verify it is not a strict subset of any other context
2344 selector and if it is not, it is the best alternative no matter
2345 whether the others do or don't match. */
2346 if (variant1
&& variant1
== variant2
)
2348 tree ctx1
= TREE_VALUE (TREE_VALUE (variant1
));
2349 FOR_EACH_VEC_ELT (variants
, i
, attr2
)
2351 if (attr2
== variant1
)
2353 tree ctx2
= TREE_VALUE (TREE_VALUE (attr2
));
2354 int r
= omp_context_selector_compare (ctx1
, ctx2
);
2357 /* The winner is a strict subset of ctx2, can't
2359 variant1
= NULL_TREE
;
2365 vec_free (entry
.variants
);
2366 return TREE_PURPOSE (TREE_VALUE (variant1
));
2370 static struct cgraph_node_hook_list
*node_removal_hook_holder
;
2371 if (!node_removal_hook_holder
)
2372 node_removal_hook_holder
2373 = symtab
->add_cgraph_removal_hook (omp_declare_variant_remove_hook
,
2376 if (omp_declare_variants
== NULL
)
2377 omp_declare_variants
2378 = hash_table
<omp_declare_variant_hasher
>::create_ggc (64);
2379 omp_declare_variant_base_entry
**slot
2380 = omp_declare_variants
->find_slot (&entry
, INSERT
);
2383 vec_free (entry
.variants
);
2384 return (*slot
)->node
->decl
;
2387 *slot
= ggc_cleared_alloc
<omp_declare_variant_base_entry
> ();
2388 (*slot
)->base
= entry
.base
;
2389 (*slot
)->node
= entry
.base
;
2390 (*slot
)->variants
= entry
.variants
;
2391 tree alt
= build_decl (DECL_SOURCE_LOCATION (base
), FUNCTION_DECL
,
2392 DECL_NAME (base
), TREE_TYPE (base
));
2393 DECL_ARTIFICIAL (alt
) = 1;
2394 DECL_IGNORED_P (alt
) = 1;
2395 TREE_STATIC (alt
) = 1;
2396 tree attributes
= DECL_ATTRIBUTES (base
);
2397 if (lookup_attribute ("noipa", attributes
) == NULL
)
2399 attributes
= tree_cons (get_identifier ("noipa"), NULL
, attributes
);
2400 if (lookup_attribute ("noinline", attributes
) == NULL
)
2401 attributes
= tree_cons (get_identifier ("noinline"), NULL
,
2403 if (lookup_attribute ("noclone", attributes
) == NULL
)
2404 attributes
= tree_cons (get_identifier ("noclone"), NULL
,
2406 if (lookup_attribute ("no_icf", attributes
) == NULL
)
2407 attributes
= tree_cons (get_identifier ("no_icf"), NULL
,
2410 DECL_ATTRIBUTES (alt
) = attributes
;
2411 DECL_INITIAL (alt
) = error_mark_node
;
2412 (*slot
)->node
= cgraph_node::create (alt
);
2413 (*slot
)->node
->declare_variant_alt
= 1;
2414 (*slot
)->node
->create_reference (entry
.base
, IPA_REF_ADDR
);
2415 omp_declare_variant_entry
*varentry
;
2416 FOR_EACH_VEC_SAFE_ELT (entry
.variants
, i
, varentry
)
2417 (*slot
)->node
->create_reference (varentry
->variant
, IPA_REF_ADDR
);
2418 if (omp_declare_variant_alt
== NULL
)
2419 omp_declare_variant_alt
2420 = hash_table
<omp_declare_variant_alt_hasher
>::create_ggc (64);
2421 *omp_declare_variant_alt
->find_slot_with_hash (*slot
, DECL_UID (alt
),
2426 if (variants
.length () == 1)
2427 return TREE_PURPOSE (TREE_VALUE (variants
[0]));
2429 /* A context selector that is a strict subset of another context selector
2430 has a score of zero. */
2433 FOR_EACH_VEC_ELT (variants
, i
, attr1
)
2436 tree ctx1
= TREE_VALUE (TREE_VALUE (attr1
));
2437 FOR_EACH_VEC_ELT_FROM (variants
, j
, attr2
, i
+ 1)
2440 tree ctx2
= TREE_VALUE (TREE_VALUE (attr2
));
2441 int r
= omp_context_selector_compare (ctx1
, ctx2
);
2444 /* ctx1 is a strict subset of ctx2, remove
2445 attr1 from the vector. */
2446 variants
[i
] = NULL_TREE
;
2450 /* ctx2 is a strict subset of ctx1, remove attr2
2452 variants
[j
] = NULL_TREE
;
2455 widest_int max_score1
= 0;
2456 widest_int max_score2
= 0;
2458 FOR_EACH_VEC_ELT (variants
, i
, attr1
)
2470 ctx
= TREE_VALUE (TREE_VALUE (variant1
));
2471 need_two
= omp_context_compute_score (ctx
, &max_score1
, false);
2473 omp_context_compute_score (ctx
, &max_score2
, true);
2475 max_score2
= max_score1
;
2477 ctx
= TREE_VALUE (TREE_VALUE (attr1
));
2478 need_two
= omp_context_compute_score (ctx
, &score1
, false);
2480 omp_context_compute_score (ctx
, &score2
, true);
2483 if (score1
> max_score1
)
2485 max_score1
= score1
;
2488 if (score2
> max_score2
)
2490 max_score2
= score2
;
2500 /* If there is a disagreement on which variant has the highest score
2501 depending on whether it will be in a declare simd clone or not,
2502 punt for now and defer until after IPA where we will know that. */
2503 return ((variant1
&& variant1
== variant2
)
2504 ? TREE_PURPOSE (TREE_VALUE (variant1
)) : base
);
2508 omp_lto_output_declare_variant_alt (lto_simple_output_block
*ob
,
2510 lto_symtab_encoder_t encoder
)
2512 gcc_assert (node
->declare_variant_alt
);
2514 omp_declare_variant_base_entry entry
;
2517 entry
.variants
= NULL
;
2518 omp_declare_variant_base_entry
*entryp
2519 = omp_declare_variant_alt
->find_with_hash (&entry
, DECL_UID (node
->decl
));
2520 gcc_assert (entryp
);
2522 int nbase
= lto_symtab_encoder_lookup (encoder
, entryp
->base
);
2523 gcc_assert (nbase
!= LCC_NOT_FOUND
);
2524 streamer_write_hwi_stream (ob
->main_stream
, nbase
);
2526 streamer_write_hwi_stream (ob
->main_stream
, entryp
->variants
->length ());
2529 omp_declare_variant_entry
*varentry
;
2530 FOR_EACH_VEC_SAFE_ELT (entryp
->variants
, i
, varentry
)
2532 int nvar
= lto_symtab_encoder_lookup (encoder
, varentry
->variant
);
2533 gcc_assert (nvar
!= LCC_NOT_FOUND
);
2534 streamer_write_hwi_stream (ob
->main_stream
, nvar
);
2536 for (widest_int
*w
= &varentry
->score
; ;
2537 w
= &varentry
->score_in_declare_simd_clone
)
2539 unsigned len
= w
->get_len ();
2540 streamer_write_hwi_stream (ob
->main_stream
, len
);
2541 const HOST_WIDE_INT
*val
= w
->get_val ();
2542 for (unsigned j
= 0; j
< len
; j
++)
2543 streamer_write_hwi_stream (ob
->main_stream
, val
[j
]);
2544 if (w
== &varentry
->score_in_declare_simd_clone
)
2548 HOST_WIDE_INT cnt
= -1;
2549 HOST_WIDE_INT i
= varentry
->matches
? 1 : 0;
2550 for (tree attr
= DECL_ATTRIBUTES (entryp
->base
->decl
);
2551 attr
; attr
= TREE_CHAIN (attr
), i
+= 2)
2553 attr
= lookup_attribute ("omp declare variant base", attr
);
2554 if (attr
== NULL_TREE
)
2557 if (varentry
->ctx
== TREE_VALUE (TREE_VALUE (attr
)))
2564 gcc_assert (cnt
!= -1);
2565 streamer_write_hwi_stream (ob
->main_stream
, cnt
);
2570 omp_lto_input_declare_variant_alt (lto_input_block
*ib
, cgraph_node
*node
,
2571 vec
<symtab_node
*> nodes
)
2573 gcc_assert (node
->declare_variant_alt
);
2574 omp_declare_variant_base_entry
*entryp
2575 = ggc_cleared_alloc
<omp_declare_variant_base_entry
> ();
2576 entryp
->base
= dyn_cast
<cgraph_node
*> (nodes
[streamer_read_hwi (ib
)]);
2577 entryp
->node
= node
;
2578 unsigned int len
= streamer_read_hwi (ib
);
2579 vec_alloc (entryp
->variants
, len
);
2581 for (unsigned int i
= 0; i
< len
; i
++)
2583 omp_declare_variant_entry varentry
;
2585 = dyn_cast
<cgraph_node
*> (nodes
[streamer_read_hwi (ib
)]);
2586 for (widest_int
*w
= &varentry
.score
; ;
2587 w
= &varentry
.score_in_declare_simd_clone
)
2589 unsigned len2
= streamer_read_hwi (ib
);
2590 HOST_WIDE_INT arr
[WIDE_INT_MAX_ELTS
];
2591 gcc_assert (len2
<= WIDE_INT_MAX_ELTS
);
2592 for (unsigned int j
= 0; j
< len2
; j
++)
2593 arr
[j
] = streamer_read_hwi (ib
);
2594 *w
= widest_int::from_array (arr
, len2
, true);
2595 if (w
== &varentry
.score_in_declare_simd_clone
)
2599 HOST_WIDE_INT cnt
= streamer_read_hwi (ib
);
2600 HOST_WIDE_INT j
= 0;
2601 varentry
.ctx
= NULL_TREE
;
2602 varentry
.matches
= (cnt
& 1) ? true : false;
2603 cnt
&= ~HOST_WIDE_INT_1
;
2604 for (tree attr
= DECL_ATTRIBUTES (entryp
->base
->decl
);
2605 attr
; attr
= TREE_CHAIN (attr
), j
+= 2)
2607 attr
= lookup_attribute ("omp declare variant base", attr
);
2608 if (attr
== NULL_TREE
)
2613 varentry
.ctx
= TREE_VALUE (TREE_VALUE (attr
));
2617 gcc_assert (varentry
.ctx
!= NULL_TREE
);
2618 entryp
->variants
->quick_push (varentry
);
2620 if (omp_declare_variant_alt
== NULL
)
2621 omp_declare_variant_alt
2622 = hash_table
<omp_declare_variant_alt_hasher
>::create_ggc (64);
2623 *omp_declare_variant_alt
->find_slot_with_hash (entryp
, DECL_UID (node
->decl
),
2627 /* Encode an oacc launch argument. This matches the GOMP_LAUNCH_PACK
2628 macro on gomp-constants.h. We do not check for overflow. */
2631 oacc_launch_pack (unsigned code
, tree device
, unsigned op
)
2635 res
= build_int_cst (unsigned_type_node
, GOMP_LAUNCH_PACK (code
, 0, op
));
2638 device
= fold_build2 (LSHIFT_EXPR
, unsigned_type_node
,
2639 device
, build_int_cst (unsigned_type_node
,
2640 GOMP_LAUNCH_DEVICE_SHIFT
));
2641 res
= fold_build2 (BIT_IOR_EXPR
, unsigned_type_node
, res
, device
);
2646 /* FIXME: What is the following comment for? */
2647 /* Look for compute grid dimension clauses and convert to an attribute
2648 attached to FN. This permits the target-side code to (a) massage
2649 the dimensions, (b) emit that data and (c) optimize. Non-constant
2650 dimensions are pushed onto ARGS.
2652 The attribute value is a TREE_LIST. A set of dimensions is
2653 represented as a list of INTEGER_CST. Those that are runtime
2654 exprs are represented as an INTEGER_CST of zero.
2656 TODO: Normally the attribute will just contain a single such list. If
2657 however it contains a list of lists, this will represent the use of
2658 device_type. Each member of the outer list is an assoc list of
2659 dimensions, keyed by the device type. The first entry will be the
2660 default. Well, that's the plan. */
2662 /* Replace any existing oacc fn attribute with updated dimensions. */
2664 /* Variant working on a list of attributes. */
2667 oacc_replace_fn_attrib_attr (tree attribs
, tree dims
)
2669 tree ident
= get_identifier (OACC_FN_ATTRIB
);
2671 /* If we happen to be present as the first attrib, drop it. */
2672 if (attribs
&& TREE_PURPOSE (attribs
) == ident
)
2673 attribs
= TREE_CHAIN (attribs
);
2674 return tree_cons (ident
, dims
, attribs
);
2677 /* Variant working on a function decl. */
2680 oacc_replace_fn_attrib (tree fn
, tree dims
)
2682 DECL_ATTRIBUTES (fn
)
2683 = oacc_replace_fn_attrib_attr (DECL_ATTRIBUTES (fn
), dims
);
2686 /* Scan CLAUSES for launch dimensions and attach them to the oacc
2687 function attribute. Push any that are non-constant onto the ARGS
2688 list, along with an appropriate GOMP_LAUNCH_DIM tag. */
2691 oacc_set_fn_attrib (tree fn
, tree clauses
, vec
<tree
> *args
)
2693 /* Must match GOMP_DIM ordering. */
2694 static const omp_clause_code ids
[]
2695 = { OMP_CLAUSE_NUM_GANGS
, OMP_CLAUSE_NUM_WORKERS
,
2696 OMP_CLAUSE_VECTOR_LENGTH
};
2698 tree dims
[GOMP_DIM_MAX
];
2700 tree attr
= NULL_TREE
;
2701 unsigned non_const
= 0;
2703 for (ix
= GOMP_DIM_MAX
; ix
--;)
2705 tree clause
= omp_find_clause (clauses
, ids
[ix
]);
2706 tree dim
= NULL_TREE
;
2709 dim
= OMP_CLAUSE_EXPR (clause
, ids
[ix
]);
2711 if (dim
&& TREE_CODE (dim
) != INTEGER_CST
)
2713 dim
= integer_zero_node
;
2714 non_const
|= GOMP_DIM_MASK (ix
);
2716 attr
= tree_cons (NULL_TREE
, dim
, attr
);
2719 oacc_replace_fn_attrib (fn
, attr
);
2723 /* Push a dynamic argument set. */
2724 args
->safe_push (oacc_launch_pack (GOMP_LAUNCH_DIM
,
2725 NULL_TREE
, non_const
));
2726 for (unsigned ix
= 0; ix
!= GOMP_DIM_MAX
; ix
++)
2727 if (non_const
& GOMP_DIM_MASK (ix
))
2728 args
->safe_push (dims
[ix
]);
2732 /* Verify OpenACC routine clauses.
2734 Returns 0 if FNDECL should be marked with an OpenACC 'routine' directive, 1
2735 if it has already been marked in compatible way, and -1 if incompatible.
2736 Upon returning, the chain of clauses will contain exactly one clause
2737 specifying the level of parallelism. */
2740 oacc_verify_routine_clauses (tree fndecl
, tree
*clauses
, location_t loc
,
2741 const char *routine_str
)
2743 tree c_level
= NULL_TREE
;
2744 tree c_nohost
= NULL_TREE
;
2745 tree c_p
= NULL_TREE
;
2746 for (tree c
= *clauses
; c
; c_p
= c
, c
= OMP_CLAUSE_CHAIN (c
))
2747 switch (OMP_CLAUSE_CODE (c
))
2749 case OMP_CLAUSE_GANG
:
2750 case OMP_CLAUSE_WORKER
:
2751 case OMP_CLAUSE_VECTOR
:
2752 case OMP_CLAUSE_SEQ
:
2753 if (c_level
== NULL_TREE
)
2755 else if (OMP_CLAUSE_CODE (c
) == OMP_CLAUSE_CODE (c_level
))
2757 /* This has already been diagnosed in the front ends. */
2758 /* Drop the duplicate clause. */
2759 gcc_checking_assert (c_p
!= NULL_TREE
);
2760 OMP_CLAUSE_CHAIN (c_p
) = OMP_CLAUSE_CHAIN (c
);
2765 error_at (OMP_CLAUSE_LOCATION (c
),
2766 "%qs specifies a conflicting level of parallelism",
2767 omp_clause_code_name
[OMP_CLAUSE_CODE (c
)]);
2768 inform (OMP_CLAUSE_LOCATION (c_level
),
2769 "... to the previous %qs clause here",
2770 omp_clause_code_name
[OMP_CLAUSE_CODE (c_level
)]);
2771 /* Drop the conflicting clause. */
2772 gcc_checking_assert (c_p
!= NULL_TREE
);
2773 OMP_CLAUSE_CHAIN (c_p
) = OMP_CLAUSE_CHAIN (c
);
2777 case OMP_CLAUSE_NOHOST
:
2778 /* Don't worry about duplicate clauses here. */
2784 if (c_level
== NULL_TREE
)
2786 /* Default to an implicit 'seq' clause. */
2787 c_level
= build_omp_clause (loc
, OMP_CLAUSE_SEQ
);
2788 OMP_CLAUSE_CHAIN (c_level
) = *clauses
;
2791 /* In *clauses, we now have exactly one clause specifying the level of
2795 = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (fndecl
));
2796 if (attr
!= NULL_TREE
)
2798 /* Diagnose if "#pragma omp declare target" has also been applied. */
2799 if (TREE_VALUE (attr
) == NULL_TREE
)
2801 /* See <https://gcc.gnu.org/PR93465>; the semantics of combining
2802 OpenACC and OpenMP 'target' are not clear. */
2804 "cannot apply %<%s%> to %qD, which has also been"
2805 " marked with an OpenMP 'declare target' directive",
2806 routine_str
, fndecl
);
2811 /* If a "#pragma acc routine" has already been applied, just verify
2812 this one for compatibility. */
2813 /* Collect previous directive's clauses. */
2814 tree c_level_p
= NULL_TREE
;
2815 tree c_nohost_p
= NULL_TREE
;
2816 for (tree c
= TREE_VALUE (attr
); c
; c
= OMP_CLAUSE_CHAIN (c
))
2817 switch (OMP_CLAUSE_CODE (c
))
2819 case OMP_CLAUSE_GANG
:
2820 case OMP_CLAUSE_WORKER
:
2821 case OMP_CLAUSE_VECTOR
:
2822 case OMP_CLAUSE_SEQ
:
2823 gcc_checking_assert (c_level_p
== NULL_TREE
);
2826 case OMP_CLAUSE_NOHOST
:
2827 gcc_checking_assert (c_nohost_p
== NULL_TREE
);
2833 gcc_checking_assert (c_level_p
!= NULL_TREE
);
2834 /* ..., and compare to current directive's, which we've already collected
2838 /* Matching level of parallelism? */
2839 if (OMP_CLAUSE_CODE (c_level
) != OMP_CLAUSE_CODE (c_level_p
))
2842 c_diag_p
= c_level_p
;
2845 /* Matching 'nohost' clauses? */
2846 if ((c_nohost
== NULL_TREE
) != (c_nohost_p
== NULL_TREE
))
2849 c_diag_p
= c_nohost_p
;
2856 if (c_diag
!= NULL_TREE
)
2857 error_at (OMP_CLAUSE_LOCATION (c_diag
),
2858 "incompatible %qs clause when applying"
2859 " %<%s%> to %qD, which has already been"
2860 " marked with an OpenACC 'routine' directive",
2861 omp_clause_code_name
[OMP_CLAUSE_CODE (c_diag
)],
2862 routine_str
, fndecl
);
2863 else if (c_diag_p
!= NULL_TREE
)
2865 "missing %qs clause when applying"
2866 " %<%s%> to %qD, which has already been"
2867 " marked with an OpenACC 'routine' directive",
2868 omp_clause_code_name
[OMP_CLAUSE_CODE (c_diag_p
)],
2869 routine_str
, fndecl
);
2872 if (c_diag_p
!= NULL_TREE
)
2873 inform (OMP_CLAUSE_LOCATION (c_diag_p
),
2874 "... with %qs clause here",
2875 omp_clause_code_name
[OMP_CLAUSE_CODE (c_diag_p
)]);
2878 /* In the front ends, we don't preserve location information for the
2879 OpenACC routine directive itself. However, that of c_level_p
2881 location_t loc_routine
= OMP_CLAUSE_LOCATION (c_level_p
);
2882 inform (loc_routine
, "... without %qs clause near to here",
2883 omp_clause_code_name
[OMP_CLAUSE_CODE (c_diag
)]);
2892 /* Process the OpenACC 'routine' directive clauses to generate an attribute
2893 for the level of parallelism. All dimensions have a size of zero
2894 (dynamic). TREE_PURPOSE is set to indicate whether that dimension
2895 can have a loop partitioned on it. non-zero indicates
2896 yes, zero indicates no. By construction once a non-zero has been
2897 reached, further inner dimensions must also be non-zero. We set
2898 TREE_VALUE to zero for the dimensions that may be partitioned and
2899 1 for the other ones -- if a loop is (erroneously) spawned at
2900 an outer level, we don't want to try and partition it. */
2903 oacc_build_routine_dims (tree clauses
)
2905 /* Must match GOMP_DIM ordering. */
2906 static const omp_clause_code ids
[]
2907 = {OMP_CLAUSE_GANG
, OMP_CLAUSE_WORKER
, OMP_CLAUSE_VECTOR
, OMP_CLAUSE_SEQ
};
2911 for (; clauses
; clauses
= OMP_CLAUSE_CHAIN (clauses
))
2912 for (ix
= GOMP_DIM_MAX
+ 1; ix
--;)
2913 if (OMP_CLAUSE_CODE (clauses
) == ids
[ix
])
2918 gcc_checking_assert (level
>= 0);
2920 tree dims
= NULL_TREE
;
2922 for (ix
= GOMP_DIM_MAX
; ix
--;)
2923 dims
= tree_cons (build_int_cst (boolean_type_node
, ix
>= level
),
2924 build_int_cst (integer_type_node
, ix
< level
), dims
);
2929 /* Retrieve the oacc function attrib and return it. Non-oacc
2930 functions will return NULL. */
2933 oacc_get_fn_attrib (tree fn
)
2935 return lookup_attribute (OACC_FN_ATTRIB
, DECL_ATTRIBUTES (fn
));
2938 /* Return true if FN is an OpenMP or OpenACC offloading function. */
2941 offloading_function_p (tree fn
)
2943 tree attrs
= DECL_ATTRIBUTES (fn
);
2944 return (lookup_attribute ("omp declare target", attrs
)
2945 || lookup_attribute ("omp target entrypoint", attrs
));
2948 /* Extract an oacc execution dimension from FN. FN must be an
2949 offloaded function or routine that has already had its execution
2950 dimensions lowered to the target-specific values. */
2953 oacc_get_fn_dim_size (tree fn
, int axis
)
2955 tree attrs
= oacc_get_fn_attrib (fn
);
2957 gcc_assert (axis
< GOMP_DIM_MAX
);
2959 tree dims
= TREE_VALUE (attrs
);
2961 dims
= TREE_CHAIN (dims
);
2963 int size
= TREE_INT_CST_LOW (TREE_VALUE (dims
));
2968 /* Extract the dimension axis from an IFN_GOACC_DIM_POS or
2969 IFN_GOACC_DIM_SIZE call. */
2972 oacc_get_ifn_dim_arg (const gimple
*stmt
)
2974 gcc_checking_assert (gimple_call_internal_fn (stmt
) == IFN_GOACC_DIM_SIZE
2975 || gimple_call_internal_fn (stmt
) == IFN_GOACC_DIM_POS
);
2976 tree arg
= gimple_call_arg (stmt
, 0);
2977 HOST_WIDE_INT axis
= TREE_INT_CST_LOW (arg
);
2979 gcc_checking_assert (axis
>= 0 && axis
< GOMP_DIM_MAX
);
2983 #include "gt-omp-general.h"