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-2020 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 "hsa-common.h"
43 #include "tree-pass.h"
44 #include "omp-device-properties.h"
46 enum omp_requires omp_requires_mask
;
49 omp_find_clause (tree clauses
, enum omp_clause_code kind
)
51 for (; clauses
; clauses
= OMP_CLAUSE_CHAIN (clauses
))
52 if (OMP_CLAUSE_CODE (clauses
) == kind
)
58 /* True if OpenMP should regard this DECL as being a scalar which has Fortran's
59 allocatable or pointer attribute. */
61 omp_is_allocatable_or_ptr (tree decl
)
63 return lang_hooks
.decls
.omp_is_allocatable_or_ptr (decl
);
66 /* Check whether this DECL belongs to a Fortran optional argument.
67 With 'for_present_check' set to false, decls which are optional parameters
68 themselve are returned as tree - or a NULL_TREE otherwise. Those decls are
69 always pointers. With 'for_present_check' set to true, the decl for checking
70 whether an argument is present is returned; for arguments with value
71 attribute this is the hidden argument and of BOOLEAN_TYPE. If the decl is
72 unrelated to optional arguments, NULL_TREE is returned. */
75 omp_check_optional_argument (tree decl
, bool for_present_check
)
77 return lang_hooks
.decls
.omp_check_optional_argument (decl
, for_present_check
);
80 /* Return true if DECL is a reference type. */
83 omp_is_reference (tree decl
)
85 return lang_hooks
.decls
.omp_privatize_by_reference (decl
);
88 /* Adjust *COND_CODE and *N2 so that the former is either LT_EXPR or GT_EXPR,
89 given that V is the loop index variable and STEP is loop step. */
92 omp_adjust_for_condition (location_t loc
, enum tree_code
*cond_code
, tree
*n2
,
102 gcc_assert (TREE_CODE (step
) == INTEGER_CST
);
103 if (TREE_CODE (TREE_TYPE (v
)) == INTEGER_TYPE
)
105 if (integer_onep (step
))
106 *cond_code
= LT_EXPR
;
109 gcc_assert (integer_minus_onep (step
));
110 *cond_code
= GT_EXPR
;
115 tree unit
= TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (v
)));
116 gcc_assert (TREE_CODE (unit
) == INTEGER_CST
);
117 if (tree_int_cst_equal (unit
, step
))
118 *cond_code
= LT_EXPR
;
121 gcc_assert (wi::neg (wi::to_widest (unit
))
122 == wi::to_widest (step
));
123 *cond_code
= GT_EXPR
;
130 if (POINTER_TYPE_P (TREE_TYPE (*n2
)))
131 *n2
= fold_build_pointer_plus_hwi_loc (loc
, *n2
, 1);
133 *n2
= fold_build2_loc (loc
, PLUS_EXPR
, TREE_TYPE (*n2
), *n2
,
134 build_int_cst (TREE_TYPE (*n2
), 1));
135 *cond_code
= LT_EXPR
;
138 if (POINTER_TYPE_P (TREE_TYPE (*n2
)))
139 *n2
= fold_build_pointer_plus_hwi_loc (loc
, *n2
, -1);
141 *n2
= fold_build2_loc (loc
, MINUS_EXPR
, TREE_TYPE (*n2
), *n2
,
142 build_int_cst (TREE_TYPE (*n2
), 1));
143 *cond_code
= GT_EXPR
;
150 /* Return the looping step from INCR, extracted from the step of a gimple omp
154 omp_get_for_step_from_incr (location_t loc
, tree incr
)
157 switch (TREE_CODE (incr
))
160 step
= TREE_OPERAND (incr
, 1);
162 case POINTER_PLUS_EXPR
:
163 step
= fold_convert (ssizetype
, TREE_OPERAND (incr
, 1));
166 step
= TREE_OPERAND (incr
, 1);
167 step
= fold_build1_loc (loc
, NEGATE_EXPR
, TREE_TYPE (step
), step
);
175 /* Extract the header elements of parallel loop FOR_STMT and store
179 omp_extract_for_data (gomp_for
*for_stmt
, struct omp_for_data
*fd
,
180 struct omp_for_data_loop
*loops
)
182 tree t
, var
, *collapse_iter
, *collapse_count
;
183 tree count
= NULL_TREE
, iter_type
= long_integer_type_node
;
184 struct omp_for_data_loop
*loop
;
186 struct omp_for_data_loop dummy_loop
;
187 location_t loc
= gimple_location (for_stmt
);
188 bool simd
= gimple_omp_for_kind (for_stmt
) == GF_OMP_FOR_KIND_SIMD
;
189 bool distribute
= gimple_omp_for_kind (for_stmt
)
190 == GF_OMP_FOR_KIND_DISTRIBUTE
;
191 bool taskloop
= gimple_omp_for_kind (for_stmt
)
192 == GF_OMP_FOR_KIND_TASKLOOP
;
195 fd
->for_stmt
= for_stmt
;
197 fd
->have_nowait
= distribute
|| simd
;
198 fd
->have_ordered
= false;
199 fd
->have_reductemp
= false;
200 fd
->have_pointer_condtemp
= false;
201 fd
->have_scantemp
= false;
202 fd
->have_nonctrl_scantemp
= false;
203 fd
->lastprivate_conditional
= 0;
204 fd
->tiling
= NULL_TREE
;
207 fd
->sched_kind
= OMP_CLAUSE_SCHEDULE_STATIC
;
208 fd
->sched_modifiers
= 0;
209 fd
->chunk_size
= NULL_TREE
;
210 fd
->simd_schedule
= false;
211 collapse_iter
= NULL
;
212 collapse_count
= NULL
;
214 for (t
= gimple_omp_for_clauses (for_stmt
); t
; t
= OMP_CLAUSE_CHAIN (t
))
215 switch (OMP_CLAUSE_CODE (t
))
217 case OMP_CLAUSE_NOWAIT
:
218 fd
->have_nowait
= true;
220 case OMP_CLAUSE_ORDERED
:
221 fd
->have_ordered
= true;
222 if (OMP_CLAUSE_ORDERED_EXPR (t
))
223 fd
->ordered
= tree_to_shwi (OMP_CLAUSE_ORDERED_EXPR (t
));
225 case OMP_CLAUSE_SCHEDULE
:
226 gcc_assert (!distribute
&& !taskloop
);
228 = (enum omp_clause_schedule_kind
)
229 (OMP_CLAUSE_SCHEDULE_KIND (t
) & OMP_CLAUSE_SCHEDULE_MASK
);
230 fd
->sched_modifiers
= (OMP_CLAUSE_SCHEDULE_KIND (t
)
231 & ~OMP_CLAUSE_SCHEDULE_MASK
);
232 fd
->chunk_size
= OMP_CLAUSE_SCHEDULE_CHUNK_EXPR (t
);
233 fd
->simd_schedule
= OMP_CLAUSE_SCHEDULE_SIMD (t
);
235 case OMP_CLAUSE_DIST_SCHEDULE
:
236 gcc_assert (distribute
);
237 fd
->chunk_size
= OMP_CLAUSE_DIST_SCHEDULE_CHUNK_EXPR (t
);
239 case OMP_CLAUSE_COLLAPSE
:
240 fd
->collapse
= tree_to_shwi (OMP_CLAUSE_COLLAPSE_EXPR (t
));
241 if (fd
->collapse
> 1)
243 collapse_iter
= &OMP_CLAUSE_COLLAPSE_ITERVAR (t
);
244 collapse_count
= &OMP_CLAUSE_COLLAPSE_COUNT (t
);
247 case OMP_CLAUSE_TILE
:
248 fd
->tiling
= OMP_CLAUSE_TILE_LIST (t
);
249 fd
->collapse
= list_length (fd
->tiling
);
250 gcc_assert (fd
->collapse
);
251 collapse_iter
= &OMP_CLAUSE_TILE_ITERVAR (t
);
252 collapse_count
= &OMP_CLAUSE_TILE_COUNT (t
);
254 case OMP_CLAUSE__REDUCTEMP_
:
255 fd
->have_reductemp
= true;
257 case OMP_CLAUSE_LASTPRIVATE
:
258 if (OMP_CLAUSE_LASTPRIVATE_CONDITIONAL (t
))
259 fd
->lastprivate_conditional
++;
261 case OMP_CLAUSE__CONDTEMP_
:
262 if (POINTER_TYPE_P (TREE_TYPE (OMP_CLAUSE_DECL (t
))))
263 fd
->have_pointer_condtemp
= true;
265 case OMP_CLAUSE__SCANTEMP_
:
266 fd
->have_scantemp
= true;
267 if (!OMP_CLAUSE__SCANTEMP__ALLOC (t
)
268 && !OMP_CLAUSE__SCANTEMP__CONTROL (t
))
269 fd
->have_nonctrl_scantemp
= true;
275 if (fd
->collapse
> 1 || fd
->tiling
)
278 fd
->loops
= &fd
->loop
;
280 if (fd
->ordered
&& fd
->collapse
== 1 && loops
!= NULL
)
285 collapse_iter
= &iterv
;
286 collapse_count
= &countv
;
289 /* FIXME: for now map schedule(auto) to schedule(static).
290 There should be analysis to determine whether all iterations
291 are approximately the same amount of work (then schedule(static)
292 is best) or if it varies (then schedule(dynamic,N) is better). */
293 if (fd
->sched_kind
== OMP_CLAUSE_SCHEDULE_AUTO
)
295 fd
->sched_kind
= OMP_CLAUSE_SCHEDULE_STATIC
;
296 gcc_assert (fd
->chunk_size
== NULL
);
298 gcc_assert ((fd
->collapse
== 1 && !fd
->tiling
) || collapse_iter
!= NULL
);
300 fd
->sched_kind
= OMP_CLAUSE_SCHEDULE_RUNTIME
;
301 if (fd
->sched_kind
== OMP_CLAUSE_SCHEDULE_RUNTIME
)
302 gcc_assert (fd
->chunk_size
== NULL
);
303 else if (fd
->chunk_size
== NULL
)
305 /* We only need to compute a default chunk size for ordered
306 static loops and dynamic loops. */
307 if (fd
->sched_kind
!= OMP_CLAUSE_SCHEDULE_STATIC
309 fd
->chunk_size
= (fd
->sched_kind
== OMP_CLAUSE_SCHEDULE_STATIC
)
310 ? integer_zero_node
: integer_one_node
;
313 int cnt
= fd
->ordered
? fd
->ordered
: fd
->collapse
;
314 for (i
= 0; i
< cnt
; i
++)
319 && (fd
->ordered
== 0 || loops
== NULL
))
321 else if (loops
!= NULL
)
326 loop
->v
= gimple_omp_for_index (for_stmt
, i
);
327 gcc_assert (SSA_VAR_P (loop
->v
));
328 gcc_assert (TREE_CODE (TREE_TYPE (loop
->v
)) == INTEGER_TYPE
329 || TREE_CODE (TREE_TYPE (loop
->v
)) == POINTER_TYPE
);
330 var
= TREE_CODE (loop
->v
) == SSA_NAME
? SSA_NAME_VAR (loop
->v
) : loop
->v
;
331 loop
->n1
= gimple_omp_for_initial (for_stmt
, i
);
333 loop
->cond_code
= gimple_omp_for_cond (for_stmt
, i
);
334 loop
->n2
= gimple_omp_for_final (for_stmt
, i
);
335 gcc_assert (loop
->cond_code
!= NE_EXPR
336 || (gimple_omp_for_kind (for_stmt
)
337 != GF_OMP_FOR_KIND_OACC_LOOP
));
339 t
= gimple_omp_for_incr (for_stmt
, i
);
340 gcc_assert (TREE_OPERAND (t
, 0) == var
);
341 loop
->step
= omp_get_for_step_from_incr (loc
, t
);
343 omp_adjust_for_condition (loc
, &loop
->cond_code
, &loop
->n2
, loop
->v
,
347 || (fd
->sched_kind
== OMP_CLAUSE_SCHEDULE_STATIC
348 && !fd
->have_ordered
))
350 if (fd
->collapse
== 1 && !fd
->tiling
)
351 iter_type
= TREE_TYPE (loop
->v
);
353 || TYPE_PRECISION (iter_type
)
354 < TYPE_PRECISION (TREE_TYPE (loop
->v
)))
356 = build_nonstandard_integer_type
357 (TYPE_PRECISION (TREE_TYPE (loop
->v
)), 1);
359 else if (iter_type
!= long_long_unsigned_type_node
)
361 if (POINTER_TYPE_P (TREE_TYPE (loop
->v
)))
362 iter_type
= long_long_unsigned_type_node
;
363 else if (TYPE_UNSIGNED (TREE_TYPE (loop
->v
))
364 && TYPE_PRECISION (TREE_TYPE (loop
->v
))
365 >= TYPE_PRECISION (iter_type
))
369 if (loop
->cond_code
== LT_EXPR
)
370 n
= fold_build2_loc (loc
, PLUS_EXPR
, TREE_TYPE (loop
->v
),
371 loop
->n2
, loop
->step
);
374 if (TREE_CODE (n
) != INTEGER_CST
375 || tree_int_cst_lt (TYPE_MAX_VALUE (iter_type
), n
))
376 iter_type
= long_long_unsigned_type_node
;
378 else if (TYPE_PRECISION (TREE_TYPE (loop
->v
))
379 > TYPE_PRECISION (iter_type
))
383 if (loop
->cond_code
== LT_EXPR
)
386 n2
= fold_build2_loc (loc
, PLUS_EXPR
, TREE_TYPE (loop
->v
),
387 loop
->n2
, loop
->step
);
391 n1
= fold_build2_loc (loc
, MINUS_EXPR
, TREE_TYPE (loop
->v
),
392 loop
->n2
, loop
->step
);
395 if (TREE_CODE (n1
) != INTEGER_CST
396 || TREE_CODE (n2
) != INTEGER_CST
397 || !tree_int_cst_lt (TYPE_MIN_VALUE (iter_type
), n1
)
398 || !tree_int_cst_lt (n2
, TYPE_MAX_VALUE (iter_type
)))
399 iter_type
= long_long_unsigned_type_node
;
403 if (i
>= fd
->collapse
)
406 if (collapse_count
&& *collapse_count
== NULL
)
408 t
= fold_binary (loop
->cond_code
, boolean_type_node
,
409 fold_convert (TREE_TYPE (loop
->v
), loop
->n1
),
410 fold_convert (TREE_TYPE (loop
->v
), loop
->n2
));
411 if (t
&& integer_zerop (t
))
412 count
= build_zero_cst (long_long_unsigned_type_node
);
413 else if ((i
== 0 || count
!= NULL_TREE
)
414 && TREE_CODE (TREE_TYPE (loop
->v
)) == INTEGER_TYPE
415 && TREE_CONSTANT (loop
->n1
)
416 && TREE_CONSTANT (loop
->n2
)
417 && TREE_CODE (loop
->step
) == INTEGER_CST
)
419 tree itype
= TREE_TYPE (loop
->v
);
421 if (POINTER_TYPE_P (itype
))
422 itype
= signed_type_for (itype
);
423 t
= build_int_cst (itype
, (loop
->cond_code
== LT_EXPR
? -1 : 1));
424 t
= fold_build2_loc (loc
, PLUS_EXPR
, itype
,
425 fold_convert_loc (loc
, itype
, loop
->step
),
427 t
= fold_build2_loc (loc
, PLUS_EXPR
, itype
, t
,
428 fold_convert_loc (loc
, itype
, loop
->n2
));
429 t
= fold_build2_loc (loc
, MINUS_EXPR
, itype
, t
,
430 fold_convert_loc (loc
, itype
, loop
->n1
));
431 if (TYPE_UNSIGNED (itype
) && loop
->cond_code
== GT_EXPR
)
433 tree step
= fold_convert_loc (loc
, itype
, loop
->step
);
434 t
= fold_build2_loc (loc
, TRUNC_DIV_EXPR
, itype
,
435 fold_build1_loc (loc
, NEGATE_EXPR
,
437 fold_build1_loc (loc
, NEGATE_EXPR
,
441 t
= fold_build2_loc (loc
, TRUNC_DIV_EXPR
, itype
, t
,
442 fold_convert_loc (loc
, itype
,
444 t
= fold_convert_loc (loc
, long_long_unsigned_type_node
, t
);
445 if (count
!= NULL_TREE
)
446 count
= fold_build2_loc (loc
, MULT_EXPR
,
447 long_long_unsigned_type_node
,
451 if (TREE_CODE (count
) != INTEGER_CST
)
454 else if (count
&& !integer_zerop (count
))
461 && (fd
->sched_kind
!= OMP_CLAUSE_SCHEDULE_STATIC
462 || fd
->have_ordered
))
464 if (!tree_int_cst_lt (count
, TYPE_MAX_VALUE (long_integer_type_node
)))
465 iter_type
= long_long_unsigned_type_node
;
467 iter_type
= long_integer_type_node
;
469 else if (collapse_iter
&& *collapse_iter
!= NULL
)
470 iter_type
= TREE_TYPE (*collapse_iter
);
471 fd
->iter_type
= iter_type
;
472 if (collapse_iter
&& *collapse_iter
== NULL
)
473 *collapse_iter
= create_tmp_var (iter_type
, ".iter");
474 if (collapse_count
&& *collapse_count
== NULL
)
477 *collapse_count
= fold_convert_loc (loc
, iter_type
, count
);
479 *collapse_count
= create_tmp_var (iter_type
, ".count");
482 if (fd
->collapse
> 1 || fd
->tiling
|| (fd
->ordered
&& loops
))
484 fd
->loop
.v
= *collapse_iter
;
485 fd
->loop
.n1
= build_int_cst (TREE_TYPE (fd
->loop
.v
), 0);
486 fd
->loop
.n2
= *collapse_count
;
487 fd
->loop
.step
= build_int_cst (TREE_TYPE (fd
->loop
.v
), 1);
488 fd
->loop
.cond_code
= LT_EXPR
;
494 /* Build a call to GOMP_barrier. */
497 omp_build_barrier (tree lhs
)
499 tree fndecl
= builtin_decl_explicit (lhs
? BUILT_IN_GOMP_BARRIER_CANCEL
500 : BUILT_IN_GOMP_BARRIER
);
501 gcall
*g
= gimple_build_call (fndecl
, 0);
503 gimple_call_set_lhs (g
, lhs
);
507 /* Return maximum possible vectorization factor for the target. */
514 || !flag_tree_loop_optimize
515 || (!flag_tree_loop_vectorize
516 && global_options_set
.x_flag_tree_loop_vectorize
))
519 auto_vector_modes modes
;
520 targetm
.vectorize
.autovectorize_vector_modes (&modes
, true);
521 if (!modes
.is_empty ())
524 for (unsigned int i
= 0; i
< modes
.length (); ++i
)
525 /* The returned modes use the smallest element size (and thus
526 the largest nunits) for the vectorization approach that they
528 vf
= ordered_max (vf
, GET_MODE_NUNITS (modes
[i
]));
532 machine_mode vqimode
= targetm
.vectorize
.preferred_simd_mode (QImode
);
533 if (GET_MODE_CLASS (vqimode
) == MODE_VECTOR_INT
)
534 return GET_MODE_NUNITS (vqimode
);
539 /* Return maximum SIMT width if offloading may target SIMT hardware. */
542 omp_max_simt_vf (void)
546 if (ENABLE_OFFLOADING
)
547 for (const char *c
= getenv ("OFFLOAD_TARGET_NAMES"); c
;)
549 if (!strncmp (c
, "nvptx", strlen ("nvptx")))
551 else if ((c
= strchr (c
, ':')))
557 /* Store the construct selectors as tree codes from last to first,
558 return their number. */
561 omp_constructor_traits_to_codes (tree ctx
, enum tree_code
*constructs
)
563 int nconstructs
= list_length (ctx
);
564 int i
= nconstructs
- 1;
565 for (tree t2
= ctx
; t2
; t2
= TREE_CHAIN (t2
), i
--)
567 const char *sel
= IDENTIFIER_POINTER (TREE_PURPOSE (t2
));
568 if (!strcmp (sel
, "target"))
569 constructs
[i
] = OMP_TARGET
;
570 else if (!strcmp (sel
, "teams"))
571 constructs
[i
] = OMP_TEAMS
;
572 else if (!strcmp (sel
, "parallel"))
573 constructs
[i
] = OMP_PARALLEL
;
574 else if (!strcmp (sel
, "for") || !strcmp (sel
, "do"))
575 constructs
[i
] = OMP_FOR
;
576 else if (!strcmp (sel
, "simd"))
577 constructs
[i
] = OMP_SIMD
;
581 gcc_assert (i
== -1);
585 /* Return true if PROP is possibly present in one of the offloading target's
586 OpenMP contexts. The format of PROPS string is always offloading target's
587 name terminated by '\0', followed by properties for that offloading
588 target separated by '\0' and terminated by another '\0'. The strings
589 are created from omp-device-properties installed files of all configured
590 offloading targets. */
593 omp_offload_device_kind_arch_isa (const char *props
, const char *prop
)
595 const char *names
= getenv ("OFFLOAD_TARGET_NAMES");
596 if (names
== NULL
|| *names
== '\0')
598 while (*props
!= '\0')
600 size_t name_len
= strlen (props
);
601 bool matches
= false;
602 for (const char *c
= names
; c
; )
604 if (strncmp (props
, c
, name_len
) == 0
605 && (c
[name_len
] == '\0'
606 || c
[name_len
] == ':'
607 || c
[name_len
] == '='))
612 else if ((c
= strchr (c
, ':')))
615 props
= props
+ name_len
+ 1;
616 while (*props
!= '\0')
618 if (matches
&& strcmp (props
, prop
) == 0)
620 props
= strchr (props
, '\0') + 1;
627 /* Return true if the current code location is or might be offloaded.
628 Return true in declare target functions, or when nested in a target
629 region or when unsure, return false otherwise. */
632 omp_maybe_offloaded (void)
634 if (!hsa_gen_requested_p ())
636 if (!ENABLE_OFFLOADING
)
638 const char *names
= getenv ("OFFLOAD_TARGET_NAMES");
639 if (names
== NULL
|| *names
== '\0')
642 if (symtab
->state
== PARSING
)
645 if (cfun
&& cfun
->after_inlining
)
647 if (current_function_decl
648 && lookup_attribute ("omp declare target",
649 DECL_ATTRIBUTES (current_function_decl
)))
651 if (cfun
&& (cfun
->curr_properties
& PROP_gimple_any
) == 0)
653 enum tree_code construct
= OMP_TARGET
;
654 if (omp_construct_selector_matches (&construct
, 1, NULL
))
660 /* Return a name from PROP, a property in selectors accepting
664 omp_context_name_list_prop (tree prop
)
666 if (TREE_PURPOSE (prop
))
667 return IDENTIFIER_POINTER (TREE_PURPOSE (prop
));
670 const char *ret
= TREE_STRING_POINTER (TREE_VALUE (prop
));
671 if ((size_t) TREE_STRING_LENGTH (TREE_VALUE (prop
)) == strlen (ret
) + 1)
677 /* Return 1 if context selector matches the current OpenMP context, 0
678 if it does not and -1 if it is unknown and need to be determined later.
679 Some properties can be checked right away during parsing (this routine),
680 others need to wait until the whole TU is parsed, others need to wait until
681 IPA, others until vectorization. */
684 omp_context_selector_matches (tree ctx
)
687 for (tree t1
= ctx
; t1
; t1
= TREE_CHAIN (t1
))
689 char set
= IDENTIFIER_POINTER (TREE_PURPOSE (t1
))[0];
692 /* For now, ignore the construct set. While something can be
693 determined already during parsing, we don't know until end of TU
694 whether additional constructs aren't added through declare variant
695 unless "omp declare variant variant" attribute exists already
696 (so in most of the cases), and we'd need to maintain set of
697 surrounding OpenMP constructs, which is better handled during
699 if (symtab
->state
== PARSING
)
705 enum tree_code constructs
[5];
707 = omp_constructor_traits_to_codes (TREE_VALUE (t1
), constructs
);
709 if (cfun
&& (cfun
->curr_properties
& PROP_gimple_any
) != 0)
711 if (!cfun
->after_inlining
)
717 for (i
= 0; i
< nconstructs
; ++i
)
718 if (constructs
[i
] == OMP_SIMD
)
725 /* If there is no simd, assume it is ok after IPA,
726 constructs should have been checked before. */
730 int r
= omp_construct_selector_matches (constructs
, nconstructs
,
738 for (tree t2
= TREE_VALUE (t1
); t2
; t2
= TREE_CHAIN (t2
))
740 const char *sel
= IDENTIFIER_POINTER (TREE_PURPOSE (t2
));
744 if (set
== 'i' && !strcmp (sel
, "vendor"))
745 for (tree t3
= TREE_VALUE (t2
); t3
; t3
= TREE_CHAIN (t3
))
747 const char *prop
= omp_context_name_list_prop (t3
);
750 if ((!strcmp (prop
, " score") && TREE_PURPOSE (t3
))
751 || !strcmp (prop
, "gnu"))
757 if (set
== 'i' && !strcmp (sel
, "extension"))
758 /* We don't support any extensions right now. */
762 if (set
== 'i' && !strcmp (sel
, "atomic_default_mem_order"))
764 if (cfun
&& (cfun
->curr_properties
& PROP_gimple_any
) != 0)
767 enum omp_memory_order omo
768 = ((enum omp_memory_order
)
770 & OMP_REQUIRES_ATOMIC_DEFAULT_MEM_ORDER
));
771 if (omo
== OMP_MEMORY_ORDER_UNSPECIFIED
)
773 /* We don't know yet, until end of TU. */
774 if (symtab
->state
== PARSING
)
780 omo
= OMP_MEMORY_ORDER_RELAXED
;
782 tree t3
= TREE_VALUE (t2
);
783 const char *prop
= IDENTIFIER_POINTER (TREE_PURPOSE (t3
));
784 if (!strcmp (prop
, " score"))
786 t3
= TREE_CHAIN (t3
);
787 prop
= IDENTIFIER_POINTER (TREE_PURPOSE (t3
));
789 if (!strcmp (prop
, "relaxed")
790 && omo
!= OMP_MEMORY_ORDER_RELAXED
)
792 else if (!strcmp (prop
, "seq_cst")
793 && omo
!= OMP_MEMORY_ORDER_SEQ_CST
)
795 else if (!strcmp (prop
, "acq_rel")
796 && omo
!= OMP_MEMORY_ORDER_ACQ_REL
)
799 if (set
== 'd' && !strcmp (sel
, "arch"))
800 for (tree t3
= TREE_VALUE (t2
); t3
; t3
= TREE_CHAIN (t3
))
802 const char *arch
= omp_context_name_list_prop (t3
);
806 if (targetm
.omp
.device_kind_arch_isa
!= NULL
)
807 r
= targetm
.omp
.device_kind_arch_isa (omp_device_arch
,
809 if (r
== 0 || (r
== -1 && symtab
->state
!= PARSING
))
811 /* If we are or might be in a target region or
812 declare target function, need to take into account
813 also offloading values. */
814 if (!omp_maybe_offloaded ())
816 if (strcmp (arch
, "hsa") == 0
817 && hsa_gen_requested_p ())
822 if (ENABLE_OFFLOADING
)
824 const char *arches
= omp_offload_device_arch
;
825 if (omp_offload_device_kind_arch_isa (arches
,
836 /* If arch matches on the host, it still might not match
837 in the offloading region. */
838 else if (omp_maybe_offloaded ())
843 if (set
== 'i' && !strcmp (sel
, "unified_address"))
845 if (cfun
&& (cfun
->curr_properties
& PROP_gimple_any
) != 0)
848 if ((omp_requires_mask
& OMP_REQUIRES_UNIFIED_ADDRESS
) == 0)
850 if (symtab
->state
== PARSING
)
857 if (set
== 'i' && !strcmp (sel
, "unified_shared_memory"))
859 if (cfun
&& (cfun
->curr_properties
& PROP_gimple_any
) != 0)
862 if ((omp_requires_mask
863 & OMP_REQUIRES_UNIFIED_SHARED_MEMORY
) == 0)
865 if (symtab
->state
== PARSING
)
874 if (set
== 'i' && !strcmp (sel
, "dynamic_allocators"))
876 if (cfun
&& (cfun
->curr_properties
& PROP_gimple_any
) != 0)
879 if ((omp_requires_mask
880 & OMP_REQUIRES_DYNAMIC_ALLOCATORS
) == 0)
882 if (symtab
->state
== PARSING
)
891 if (set
== 'i' && !strcmp (sel
, "reverse_offload"))
893 if (cfun
&& (cfun
->curr_properties
& PROP_gimple_any
) != 0)
896 if ((omp_requires_mask
& OMP_REQUIRES_REVERSE_OFFLOAD
) == 0)
898 if (symtab
->state
== PARSING
)
907 if (set
== 'd' && !strcmp (sel
, "kind"))
908 for (tree t3
= TREE_VALUE (t2
); t3
; t3
= TREE_CHAIN (t3
))
910 const char *prop
= omp_context_name_list_prop (t3
);
913 if (!strcmp (prop
, "any"))
915 if (!strcmp (prop
, "host"))
917 if (omp_maybe_offloaded ())
921 if (!strcmp (prop
, "nohost"))
923 if (omp_maybe_offloaded ())
930 if (targetm
.omp
.device_kind_arch_isa
!= NULL
)
931 r
= targetm
.omp
.device_kind_arch_isa (omp_device_kind
,
934 r
= strcmp (prop
, "cpu") == 0;
935 if (r
== 0 || (r
== -1 && symtab
->state
!= PARSING
))
937 /* If we are or might be in a target region or
938 declare target function, need to take into account
939 also offloading values. */
940 if (!omp_maybe_offloaded ())
942 if (strcmp (prop
, "gpu") == 0
943 && hsa_gen_requested_p ())
948 if (ENABLE_OFFLOADING
)
950 const char *kinds
= omp_offload_device_kind
;
951 if (omp_offload_device_kind_arch_isa (kinds
, prop
))
961 /* If kind matches on the host, it still might not match
962 in the offloading region. */
963 else if (omp_maybe_offloaded ())
968 if (set
== 'd' && !strcmp (sel
, "isa"))
969 for (tree t3
= TREE_VALUE (t2
); t3
; t3
= TREE_CHAIN (t3
))
971 const char *isa
= omp_context_name_list_prop (t3
);
975 if (targetm
.omp
.device_kind_arch_isa
!= NULL
)
976 r
= targetm
.omp
.device_kind_arch_isa (omp_device_isa
,
978 if (r
== 0 || (r
== -1 && symtab
->state
!= PARSING
))
980 /* If isa is valid on the target, but not in the
981 current function and current function has
982 #pragma omp declare simd on it, some simd clones
983 might have the isa added later on. */
985 && targetm
.simd_clone
.compute_vecsize_and_simdlen
986 && (cfun
== NULL
|| !cfun
->after_inlining
))
989 = DECL_ATTRIBUTES (current_function_decl
);
990 if (lookup_attribute ("omp declare simd", attrs
))
996 /* If we are or might be in a target region or
997 declare target function, need to take into account
998 also offloading values. */
999 if (!omp_maybe_offloaded ())
1001 if (ENABLE_OFFLOADING
)
1003 const char *isas
= omp_offload_device_isa
;
1004 if (omp_offload_device_kind_arch_isa (isas
, isa
))
1014 /* If isa matches on the host, it still might not match
1015 in the offloading region. */
1016 else if (omp_maybe_offloaded ())
1021 if (set
== 'u' && !strcmp (sel
, "condition"))
1022 for (tree t3
= TREE_VALUE (t2
); t3
; t3
= TREE_CHAIN (t3
))
1023 if (TREE_PURPOSE (t3
) == NULL_TREE
)
1025 if (integer_zerop (TREE_VALUE (t3
)))
1027 if (integer_nonzerop (TREE_VALUE (t3
)))
1040 /* Compare construct={simd} CLAUSES1 with CLAUSES2, return 0/-1/1/2 as
1041 in omp_context_selector_set_compare. */
1044 omp_construct_simd_compare (tree clauses1
, tree clauses2
)
1046 if (clauses1
== NULL_TREE
)
1047 return clauses2
== NULL_TREE
? 0 : -1;
1048 if (clauses2
== NULL_TREE
)
1052 struct declare_variant_simd_data
{
1053 bool inbranch
, notinbranch
;
1055 auto_vec
<tree
,16> data_sharing
;
1056 auto_vec
<tree
,16> aligned
;
1057 declare_variant_simd_data ()
1058 : inbranch(false), notinbranch(false), simdlen(NULL_TREE
) {}
1061 for (i
= 0; i
< 2; i
++)
1062 for (tree c
= i
? clauses2
: clauses1
; c
; c
= OMP_CLAUSE_CHAIN (c
))
1065 switch (OMP_CLAUSE_CODE (c
))
1067 case OMP_CLAUSE_INBRANCH
:
1068 data
[i
].inbranch
= true;
1070 case OMP_CLAUSE_NOTINBRANCH
:
1071 data
[i
].notinbranch
= true;
1073 case OMP_CLAUSE_SIMDLEN
:
1074 data
[i
].simdlen
= OMP_CLAUSE_SIMDLEN_EXPR (c
);
1076 case OMP_CLAUSE_UNIFORM
:
1077 case OMP_CLAUSE_LINEAR
:
1078 v
= &data
[i
].data_sharing
;
1080 case OMP_CLAUSE_ALIGNED
:
1081 v
= &data
[i
].aligned
;
1086 unsigned HOST_WIDE_INT argno
= tree_to_uhwi (OMP_CLAUSE_DECL (c
));
1087 if (argno
>= v
->length ())
1088 v
->safe_grow_cleared (argno
+ 1);
1091 /* Here, r is used as a bitmask, 2 is set if CLAUSES1 has something
1092 CLAUSES2 doesn't, 1 is set if CLAUSES2 has something CLAUSES1
1093 doesn't. Thus, r == 3 implies return value 2, r == 1 implies
1094 -1, r == 2 implies 1 and r == 0 implies 0. */
1095 if (data
[0].inbranch
!= data
[1].inbranch
)
1096 r
|= data
[0].inbranch
? 2 : 1;
1097 if (data
[0].notinbranch
!= data
[1].notinbranch
)
1098 r
|= data
[0].notinbranch
? 2 : 1;
1099 if (!simple_cst_equal (data
[0].simdlen
, data
[1].simdlen
))
1101 if (data
[0].simdlen
&& data
[1].simdlen
)
1103 r
|= data
[0].simdlen
? 2 : 1;
1105 if (data
[0].data_sharing
.length () < data
[1].data_sharing
.length ()
1106 || data
[0].aligned
.length () < data
[1].aligned
.length ())
1109 FOR_EACH_VEC_ELT (data
[0].data_sharing
, i
, c1
)
1111 c2
= (i
< data
[1].data_sharing
.length ()
1112 ? data
[1].data_sharing
[i
] : NULL_TREE
);
1113 if ((c1
== NULL_TREE
) != (c2
== NULL_TREE
))
1115 r
|= c1
!= NULL_TREE
? 2 : 1;
1118 if (c1
== NULL_TREE
)
1120 if (OMP_CLAUSE_CODE (c1
) != OMP_CLAUSE_CODE (c2
))
1122 if (OMP_CLAUSE_CODE (c1
) != OMP_CLAUSE_LINEAR
)
1124 if (OMP_CLAUSE_LINEAR_VARIABLE_STRIDE (c1
)
1125 != OMP_CLAUSE_LINEAR_VARIABLE_STRIDE (c2
))
1127 if (OMP_CLAUSE_LINEAR_KIND (c1
) != OMP_CLAUSE_LINEAR_KIND (c2
))
1129 if (!simple_cst_equal (OMP_CLAUSE_LINEAR_STEP (c1
),
1130 OMP_CLAUSE_LINEAR_STEP (c2
)))
1133 FOR_EACH_VEC_ELT (data
[0].aligned
, i
, c1
)
1135 c2
= i
< data
[1].aligned
.length () ? data
[1].aligned
[i
] : NULL_TREE
;
1136 if ((c1
== NULL_TREE
) != (c2
== NULL_TREE
))
1138 r
|= c1
!= NULL_TREE
? 2 : 1;
1141 if (c1
== NULL_TREE
)
1143 if (!simple_cst_equal (OMP_CLAUSE_ALIGNED_ALIGNMENT (c1
),
1144 OMP_CLAUSE_ALIGNED_ALIGNMENT (c2
)))
1153 default: gcc_unreachable ();
1157 /* Compare properties of selectors SEL from SET other than construct.
1158 Return 0/-1/1/2 as in omp_context_selector_set_compare.
1159 Unlike set names or selector names, properties can have duplicates. */
1162 omp_context_selector_props_compare (const char *set
, const char *sel
,
1163 tree ctx1
, tree ctx2
)
1166 for (int pass
= 0; pass
< 2; pass
++)
1167 for (tree t1
= pass
? ctx2
: ctx1
; t1
; t1
= TREE_CHAIN (t1
))
1170 for (t2
= pass
? ctx1
: ctx2
; t2
; t2
= TREE_CHAIN (t2
))
1171 if (TREE_PURPOSE (t1
) == TREE_PURPOSE (t2
))
1173 if (TREE_PURPOSE (t1
) == NULL_TREE
)
1175 if (set
[0] == 'u' && strcmp (sel
, "condition") == 0)
1177 if (integer_zerop (TREE_VALUE (t1
))
1178 != integer_zerop (TREE_VALUE (t2
)))
1182 if (simple_cst_equal (TREE_VALUE (t1
), TREE_VALUE (t2
)))
1185 else if (strcmp (IDENTIFIER_POINTER (TREE_PURPOSE (t1
)),
1188 if (!simple_cst_equal (TREE_VALUE (t1
), TREE_VALUE (t2
)))
1195 else if (TREE_PURPOSE (t1
)
1196 && TREE_PURPOSE (t2
) == NULL_TREE
1197 && TREE_CODE (TREE_VALUE (t2
)) == STRING_CST
)
1199 const char *p1
= omp_context_name_list_prop (t1
);
1200 const char *p2
= omp_context_name_list_prop (t2
);
1202 && strcmp (p1
, p2
) == 0
1203 && strcmp (p1
, " score"))
1206 else if (TREE_PURPOSE (t1
) == NULL_TREE
1207 && TREE_PURPOSE (t2
)
1208 && TREE_CODE (TREE_VALUE (t1
)) == STRING_CST
)
1210 const char *p1
= omp_context_name_list_prop (t1
);
1211 const char *p2
= omp_context_name_list_prop (t2
);
1213 && strcmp (p1
, p2
) == 0
1214 && strcmp (p1
, " score"))
1217 if (t2
== NULL_TREE
)
1219 int r
= pass
? -1 : 1;
1220 if (ret
&& ret
!= r
)
1234 /* Compare single context selector sets CTX1 and CTX2 with SET name.
1235 Return 0 if CTX1 is equal to CTX2,
1236 -1 if CTX1 is a strict subset of CTX2,
1237 1 if CTX2 is a strict subset of CTX1, or
1238 2 if neither context is a subset of another one. */
1241 omp_context_selector_set_compare (const char *set
, tree ctx1
, tree ctx2
)
1243 bool swapped
= false;
1245 int len1
= list_length (ctx1
);
1246 int len2
= list_length (ctx2
);
1251 std::swap (ctx1
, ctx2
);
1252 std::swap (len1
, len2
);
1258 tree simd
= get_identifier ("simd");
1259 /* Handle construct set specially. In this case the order
1260 of the selector matters too. */
1261 for (t1
= ctx1
; t1
; t1
= TREE_CHAIN (t1
))
1262 if (TREE_PURPOSE (t1
) == TREE_PURPOSE (t2
))
1265 if (TREE_PURPOSE (t1
) == simd
)
1266 r
= omp_construct_simd_compare (TREE_VALUE (t1
),
1268 if (r
== 2 || (ret
&& r
&& (ret
< 0) != (r
< 0)))
1272 t2
= TREE_CHAIN (t2
);
1273 if (t2
== NULL_TREE
)
1275 t1
= TREE_CHAIN (t1
);
1283 if (t2
!= NULL_TREE
)
1285 if (t1
!= NULL_TREE
)
1293 return swapped
? -ret
: ret
;
1295 for (tree t1
= ctx1
; t1
; t1
= TREE_CHAIN (t1
))
1298 for (t2
= ctx2
; t2
; t2
= TREE_CHAIN (t2
))
1299 if (TREE_PURPOSE (t1
) == TREE_PURPOSE (t2
))
1301 const char *sel
= IDENTIFIER_POINTER (TREE_PURPOSE (t1
));
1302 int r
= omp_context_selector_props_compare (set
, sel
,
1305 if (r
== 2 || (ret
&& r
&& (ret
< 0) != (r
< 0)))
1312 if (t2
== NULL_TREE
)
1323 return swapped
? -ret
: ret
;
1326 /* Compare whole context selector specification CTX1 and CTX2.
1327 Return 0 if CTX1 is equal to CTX2,
1328 -1 if CTX1 is a strict subset of CTX2,
1329 1 if CTX2 is a strict subset of CTX1, or
1330 2 if neither context is a subset of another one. */
1333 omp_context_selector_compare (tree ctx1
, tree ctx2
)
1335 bool swapped
= false;
1337 int len1
= list_length (ctx1
);
1338 int len2
= list_length (ctx2
);
1343 std::swap (ctx1
, ctx2
);
1344 std::swap (len1
, len2
);
1346 for (tree t1
= ctx1
; t1
; t1
= TREE_CHAIN (t1
))
1349 for (t2
= ctx2
; t2
; t2
= TREE_CHAIN (t2
))
1350 if (TREE_PURPOSE (t1
) == TREE_PURPOSE (t2
))
1352 const char *set
= IDENTIFIER_POINTER (TREE_PURPOSE (t1
));
1353 int r
= omp_context_selector_set_compare (set
, TREE_VALUE (t1
),
1355 if (r
== 2 || (ret
&& r
&& (ret
< 0) != (r
< 0)))
1362 if (t2
== NULL_TREE
)
1373 return swapped
? -ret
: ret
;
1376 /* From context selector CTX, return trait-selector with name SEL in
1377 trait-selector-set with name SET if any, or NULL_TREE if not found.
1378 If SEL is NULL, return the list of trait-selectors in SET. */
1381 omp_get_context_selector (tree ctx
, const char *set
, const char *sel
)
1383 tree setid
= get_identifier (set
);
1384 tree selid
= sel
? get_identifier (sel
) : NULL_TREE
;
1385 for (tree t1
= ctx
; t1
; t1
= TREE_CHAIN (t1
))
1386 if (TREE_PURPOSE (t1
) == setid
)
1389 return TREE_VALUE (t1
);
1390 for (tree t2
= TREE_VALUE (t1
); t2
; t2
= TREE_CHAIN (t2
))
1391 if (TREE_PURPOSE (t2
) == selid
)
1397 /* Compute *SCORE for context selector CTX. Return true if the score
1398 would be different depending on whether it is a declare simd clone or
1399 not. DECLARE_SIMD should be true for the case when it would be
1400 a declare simd clone. */
1403 omp_context_compute_score (tree ctx
, widest_int
*score
, bool declare_simd
)
1405 tree construct
= omp_get_context_selector (ctx
, "construct", NULL
);
1406 bool has_kind
= omp_get_context_selector (ctx
, "device", "kind");
1407 bool has_arch
= omp_get_context_selector (ctx
, "device", "arch");
1408 bool has_isa
= omp_get_context_selector (ctx
, "device", "isa");
1411 for (tree t1
= ctx
; t1
; t1
= TREE_CHAIN (t1
))
1412 if (TREE_VALUE (t1
) != construct
)
1413 for (tree t2
= TREE_VALUE (t1
); t2
; t2
= TREE_CHAIN (t2
))
1414 if (tree t3
= TREE_VALUE (t2
))
1415 if (TREE_PURPOSE (t3
)
1416 && strcmp (IDENTIFIER_POINTER (TREE_PURPOSE (t3
)), " score") == 0
1417 && TREE_CODE (TREE_VALUE (t3
)) == INTEGER_CST
)
1418 *score
+= wi::to_widest (TREE_VALUE (t3
));
1419 if (construct
|| has_kind
|| has_arch
|| has_isa
)
1422 enum tree_code constructs
[5];
1423 int nconstructs
= 0;
1425 nconstructs
= omp_constructor_traits_to_codes (construct
, constructs
);
1426 if (omp_construct_selector_matches (constructs
, nconstructs
, scores
)
1429 int b
= declare_simd
? nconstructs
+ 1 : 0;
1430 if (scores
[b
+ nconstructs
] + 4U < score
->get_precision ())
1432 for (int n
= 0; n
< nconstructs
; ++n
)
1434 if (scores
[b
+ n
] < 0)
1439 *score
+= wi::shifted_mask
<widest_int
> (scores
[b
+ n
], 1, false);
1442 *score
+= wi::shifted_mask
<widest_int
> (scores
[b
+ nconstructs
],
1445 *score
+= wi::shifted_mask
<widest_int
> (scores
[b
+ nconstructs
] + 1,
1448 *score
+= wi::shifted_mask
<widest_int
> (scores
[b
+ nconstructs
] + 2,
1451 else /* FIXME: Implement this. */
1457 /* Class describing a single variant. */
1458 struct GTY(()) omp_declare_variant_entry
{
1459 /* NODE of the variant. */
1460 cgraph_node
*variant
;
1461 /* Score if not in declare simd clone. */
1463 /* Score if in declare simd clone. */
1464 widest_int score_in_declare_simd_clone
;
1465 /* Context selector for the variant. */
1467 /* True if the context selector is known to match already. */
1471 /* Class describing a function with variants. */
1472 struct GTY((for_user
)) omp_declare_variant_base_entry
{
1473 /* NODE of the base function. */
1475 /* NODE of the artificial function created for the deferred variant
1478 /* Vector of the variants. */
1479 vec
<omp_declare_variant_entry
, va_gc
> *variants
;
1482 struct omp_declare_variant_hasher
1483 : ggc_ptr_hash
<omp_declare_variant_base_entry
> {
1484 static hashval_t
hash (omp_declare_variant_base_entry
*);
1485 static bool equal (omp_declare_variant_base_entry
*,
1486 omp_declare_variant_base_entry
*);
1490 omp_declare_variant_hasher::hash (omp_declare_variant_base_entry
*x
)
1492 inchash::hash hstate
;
1493 hstate
.add_int (DECL_UID (x
->base
->decl
));
1494 hstate
.add_int (x
->variants
->length ());
1495 omp_declare_variant_entry
*variant
;
1497 FOR_EACH_VEC_SAFE_ELT (x
->variants
, i
, variant
)
1499 hstate
.add_int (DECL_UID (variant
->variant
->decl
));
1500 hstate
.add_wide_int (variant
->score
);
1501 hstate
.add_wide_int (variant
->score_in_declare_simd_clone
);
1502 hstate
.add_ptr (variant
->ctx
);
1503 hstate
.add_int (variant
->matches
);
1505 return hstate
.end ();
1509 omp_declare_variant_hasher::equal (omp_declare_variant_base_entry
*x
,
1510 omp_declare_variant_base_entry
*y
)
1512 if (x
->base
!= y
->base
1513 || x
->variants
->length () != y
->variants
->length ())
1515 omp_declare_variant_entry
*variant
;
1517 FOR_EACH_VEC_SAFE_ELT (x
->variants
, i
, variant
)
1518 if (variant
->variant
!= (*y
->variants
)[i
].variant
1519 || variant
->score
!= (*y
->variants
)[i
].score
1520 || (variant
->score_in_declare_simd_clone
1521 != (*y
->variants
)[i
].score_in_declare_simd_clone
)
1522 || variant
->ctx
!= (*y
->variants
)[i
].ctx
1523 || variant
->matches
!= (*y
->variants
)[i
].matches
)
1528 static GTY(()) hash_table
<omp_declare_variant_hasher
> *omp_declare_variants
;
1530 struct omp_declare_variant_alt_hasher
1531 : ggc_ptr_hash
<omp_declare_variant_base_entry
> {
1532 static hashval_t
hash (omp_declare_variant_base_entry
*);
1533 static bool equal (omp_declare_variant_base_entry
*,
1534 omp_declare_variant_base_entry
*);
1538 omp_declare_variant_alt_hasher::hash (omp_declare_variant_base_entry
*x
)
1540 return DECL_UID (x
->node
->decl
);
1544 omp_declare_variant_alt_hasher::equal (omp_declare_variant_base_entry
*x
,
1545 omp_declare_variant_base_entry
*y
)
1547 return x
->node
== y
->node
;
1550 static GTY(()) hash_table
<omp_declare_variant_alt_hasher
>
1551 *omp_declare_variant_alt
;
1553 /* Try to resolve declare variant after gimplification. */
1556 omp_resolve_late_declare_variant (tree alt
)
1558 cgraph_node
*node
= cgraph_node::get (alt
);
1559 cgraph_node
*cur_node
= cgraph_node::get (cfun
->decl
);
1561 || !node
->declare_variant_alt
1562 || !cfun
->after_inlining
)
1565 omp_declare_variant_base_entry entry
;
1568 entry
.variants
= NULL
;
1569 omp_declare_variant_base_entry
*entryp
1570 = omp_declare_variant_alt
->find_with_hash (&entry
, DECL_UID (alt
));
1573 omp_declare_variant_entry
*varentry1
, *varentry2
;
1574 auto_vec
<bool, 16> matches
;
1575 unsigned int nmatches
= 0;
1576 FOR_EACH_VEC_SAFE_ELT (entryp
->variants
, i
, varentry1
)
1578 if (varentry1
->matches
)
1580 /* This has been checked to be ok already. */
1581 matches
.safe_push (true);
1585 switch (omp_context_selector_matches (varentry1
->ctx
))
1588 matches
.safe_push (false);
1593 matches
.safe_push (true);
1600 return entryp
->base
->decl
;
1602 /* A context selector that is a strict subset of another context selector
1603 has a score of zero. */
1604 FOR_EACH_VEC_SAFE_ELT (entryp
->variants
, i
, varentry1
)
1608 vec_safe_iterate (entryp
->variants
, j
, &varentry2
); ++j
)
1611 int r
= omp_context_selector_compare (varentry1
->ctx
,
1615 /* ctx1 is a strict subset of ctx2, ignore ctx1. */
1620 /* ctx2 is a strict subset of ctx1, remove ctx2. */
1625 widest_int max_score
= -1;
1627 FOR_EACH_VEC_SAFE_ELT (entryp
->variants
, i
, varentry1
)
1631 = (cur_node
->simdclone
? varentry1
->score_in_declare_simd_clone
1632 : varentry1
->score
);
1633 if (score
> max_score
)
1636 varentry2
= varentry1
;
1639 return varentry2
->variant
->decl
;
1642 /* Try to resolve declare variant, return the variant decl if it should
1643 be used instead of base, or base otherwise. */
1646 omp_resolve_declare_variant (tree base
)
1648 tree variant1
= NULL_TREE
, variant2
= NULL_TREE
;
1649 if (cfun
&& (cfun
->curr_properties
& PROP_gimple_any
) != 0)
1650 return omp_resolve_late_declare_variant (base
);
1652 auto_vec
<tree
, 16> variants
;
1653 auto_vec
<bool, 16> defer
;
1654 bool any_deferred
= false;
1655 for (tree attr
= DECL_ATTRIBUTES (base
); attr
; attr
= TREE_CHAIN (attr
))
1657 attr
= lookup_attribute ("omp declare variant base", attr
);
1658 if (attr
== NULL_TREE
)
1660 if (TREE_CODE (TREE_PURPOSE (TREE_VALUE (attr
))) != FUNCTION_DECL
)
1662 switch (omp_context_selector_matches (TREE_VALUE (TREE_VALUE (attr
))))
1665 /* No match, ignore. */
1668 /* Needs to be deferred. */
1669 any_deferred
= true;
1670 variants
.safe_push (attr
);
1671 defer
.safe_push (true);
1674 variants
.safe_push (attr
);
1675 defer
.safe_push (false);
1679 if (variants
.length () == 0)
1684 widest_int max_score1
= 0;
1685 widest_int max_score2
= 0;
1689 omp_declare_variant_base_entry entry
;
1690 entry
.base
= cgraph_node::get_create (base
);
1692 vec_alloc (entry
.variants
, variants
.length ());
1693 FOR_EACH_VEC_ELT (variants
, i
, attr1
)
1698 tree ctx
= TREE_VALUE (TREE_VALUE (attr1
));
1699 need_two
= omp_context_compute_score (ctx
, &score1
, false);
1701 omp_context_compute_score (ctx
, &score2
, true);
1707 max_score1
= score1
;
1708 max_score2
= score2
;
1717 if (max_score1
== score1
)
1718 variant1
= NULL_TREE
;
1719 else if (score1
> max_score1
)
1721 max_score1
= score1
;
1722 variant1
= defer
[i
] ? NULL_TREE
: attr1
;
1724 if (max_score2
== score2
)
1725 variant2
= NULL_TREE
;
1726 else if (score2
> max_score2
)
1728 max_score2
= score2
;
1729 variant2
= defer
[i
] ? NULL_TREE
: attr1
;
1732 omp_declare_variant_entry varentry
;
1734 = cgraph_node::get_create (TREE_PURPOSE (TREE_VALUE (attr1
)));
1735 varentry
.score
= score1
;
1736 varentry
.score_in_declare_simd_clone
= score2
;
1738 varentry
.matches
= !defer
[i
];
1739 entry
.variants
->quick_push (varentry
);
1742 /* If there is a clear winner variant with the score which is not
1743 deferred, verify it is not a strict subset of any other context
1744 selector and if it is not, it is the best alternative no matter
1745 whether the others do or don't match. */
1746 if (variant1
&& variant1
== variant2
)
1748 tree ctx1
= TREE_VALUE (TREE_VALUE (variant1
));
1749 FOR_EACH_VEC_ELT (variants
, i
, attr2
)
1751 if (attr2
== variant1
)
1753 tree ctx2
= TREE_VALUE (TREE_VALUE (attr2
));
1754 int r
= omp_context_selector_compare (ctx1
, ctx2
);
1757 /* The winner is a strict subset of ctx2, can't
1759 variant1
= NULL_TREE
;
1765 vec_free (entry
.variants
);
1766 return TREE_PURPOSE (TREE_VALUE (variant1
));
1770 if (omp_declare_variants
== NULL
)
1771 omp_declare_variants
1772 = hash_table
<omp_declare_variant_hasher
>::create_ggc (64);
1773 omp_declare_variant_base_entry
**slot
1774 = omp_declare_variants
->find_slot (&entry
, INSERT
);
1777 vec_free (entry
.variants
);
1778 return (*slot
)->node
->decl
;
1781 *slot
= ggc_cleared_alloc
<omp_declare_variant_base_entry
> ();
1782 (*slot
)->base
= entry
.base
;
1783 (*slot
)->node
= entry
.base
;
1784 (*slot
)->variants
= entry
.variants
;
1785 tree alt
= build_decl (DECL_SOURCE_LOCATION (base
), FUNCTION_DECL
,
1786 DECL_NAME (base
), TREE_TYPE (base
));
1787 DECL_ARTIFICIAL (alt
) = 1;
1788 DECL_IGNORED_P (alt
) = 1;
1789 TREE_STATIC (alt
) = 1;
1790 tree attributes
= DECL_ATTRIBUTES (base
);
1791 if (lookup_attribute ("noipa", attributes
) == NULL
)
1793 attributes
= tree_cons (get_identifier ("noipa"), NULL
, attributes
);
1794 if (lookup_attribute ("noinline", attributes
) == NULL
)
1795 attributes
= tree_cons (get_identifier ("noinline"), NULL
,
1797 if (lookup_attribute ("noclone", attributes
) == NULL
)
1798 attributes
= tree_cons (get_identifier ("noclone"), NULL
,
1800 if (lookup_attribute ("no_icf", attributes
) == NULL
)
1801 attributes
= tree_cons (get_identifier ("no_icf"), NULL
,
1804 DECL_ATTRIBUTES (alt
) = attributes
;
1805 DECL_INITIAL (alt
) = error_mark_node
;
1806 (*slot
)->node
= cgraph_node::create (alt
);
1807 (*slot
)->node
->declare_variant_alt
= 1;
1808 (*slot
)->node
->create_reference (entry
.base
, IPA_REF_ADDR
);
1809 omp_declare_variant_entry
*varentry
;
1810 FOR_EACH_VEC_SAFE_ELT (entry
.variants
, i
, varentry
)
1811 (*slot
)->node
->create_reference (varentry
->variant
, IPA_REF_ADDR
);
1812 if (omp_declare_variant_alt
== NULL
)
1813 omp_declare_variant_alt
1814 = hash_table
<omp_declare_variant_alt_hasher
>::create_ggc (64);
1815 *omp_declare_variant_alt
->find_slot_with_hash (*slot
, DECL_UID (alt
),
1820 if (variants
.length () == 1)
1821 return TREE_PURPOSE (TREE_VALUE (variants
[0]));
1823 /* A context selector that is a strict subset of another context selector
1824 has a score of zero. */
1827 FOR_EACH_VEC_ELT (variants
, i
, attr1
)
1830 tree ctx1
= TREE_VALUE (TREE_VALUE (attr1
));
1831 FOR_EACH_VEC_ELT_FROM (variants
, j
, attr2
, i
+ 1)
1834 tree ctx2
= TREE_VALUE (TREE_VALUE (attr2
));
1835 int r
= omp_context_selector_compare (ctx1
, ctx2
);
1838 /* ctx1 is a strict subset of ctx2, remove
1839 attr1 from the vector. */
1840 variants
[i
] = NULL_TREE
;
1844 /* ctx2 is a strict subset of ctx1, remove attr2
1846 variants
[j
] = NULL_TREE
;
1849 widest_int max_score1
= 0;
1850 widest_int max_score2
= 0;
1852 FOR_EACH_VEC_ELT (variants
, i
, attr1
)
1864 ctx
= TREE_VALUE (TREE_VALUE (variant1
));
1865 need_two
= omp_context_compute_score (ctx
, &max_score1
, false);
1867 omp_context_compute_score (ctx
, &max_score2
, true);
1869 max_score2
= max_score1
;
1871 ctx
= TREE_VALUE (TREE_VALUE (attr1
));
1872 need_two
= omp_context_compute_score (ctx
, &score1
, false);
1874 omp_context_compute_score (ctx
, &score2
, true);
1877 if (score1
> max_score1
)
1879 max_score1
= score1
;
1882 if (score2
> max_score2
)
1884 max_score2
= score2
;
1894 /* If there is a disagreement on which variant has the highest score
1895 depending on whether it will be in a declare simd clone or not,
1896 punt for now and defer until after IPA where we will know that. */
1897 return ((variant1
&& variant1
== variant2
)
1898 ? TREE_PURPOSE (TREE_VALUE (variant1
)) : base
);
1902 /* Encode an oacc launch argument. This matches the GOMP_LAUNCH_PACK
1903 macro on gomp-constants.h. We do not check for overflow. */
1906 oacc_launch_pack (unsigned code
, tree device
, unsigned op
)
1910 res
= build_int_cst (unsigned_type_node
, GOMP_LAUNCH_PACK (code
, 0, op
));
1913 device
= fold_build2 (LSHIFT_EXPR
, unsigned_type_node
,
1914 device
, build_int_cst (unsigned_type_node
,
1915 GOMP_LAUNCH_DEVICE_SHIFT
));
1916 res
= fold_build2 (BIT_IOR_EXPR
, unsigned_type_node
, res
, device
);
1921 /* FIXME: What is the following comment for? */
1922 /* Look for compute grid dimension clauses and convert to an attribute
1923 attached to FN. This permits the target-side code to (a) massage
1924 the dimensions, (b) emit that data and (c) optimize. Non-constant
1925 dimensions are pushed onto ARGS.
1927 The attribute value is a TREE_LIST. A set of dimensions is
1928 represented as a list of INTEGER_CST. Those that are runtime
1929 exprs are represented as an INTEGER_CST of zero.
1931 TODO: Normally the attribute will just contain a single such list. If
1932 however it contains a list of lists, this will represent the use of
1933 device_type. Each member of the outer list is an assoc list of
1934 dimensions, keyed by the device type. The first entry will be the
1935 default. Well, that's the plan. */
1937 /* Replace any existing oacc fn attribute with updated dimensions. */
1939 /* Variant working on a list of attributes. */
1942 oacc_replace_fn_attrib_attr (tree attribs
, tree dims
)
1944 tree ident
= get_identifier (OACC_FN_ATTRIB
);
1946 /* If we happen to be present as the first attrib, drop it. */
1947 if (attribs
&& TREE_PURPOSE (attribs
) == ident
)
1948 attribs
= TREE_CHAIN (attribs
);
1949 return tree_cons (ident
, dims
, attribs
);
1952 /* Variant working on a function decl. */
1955 oacc_replace_fn_attrib (tree fn
, tree dims
)
1957 DECL_ATTRIBUTES (fn
)
1958 = oacc_replace_fn_attrib_attr (DECL_ATTRIBUTES (fn
), dims
);
1961 /* Scan CLAUSES for launch dimensions and attach them to the oacc
1962 function attribute. Push any that are non-constant onto the ARGS
1963 list, along with an appropriate GOMP_LAUNCH_DIM tag. */
1966 oacc_set_fn_attrib (tree fn
, tree clauses
, vec
<tree
> *args
)
1968 /* Must match GOMP_DIM ordering. */
1969 static const omp_clause_code ids
[]
1970 = { OMP_CLAUSE_NUM_GANGS
, OMP_CLAUSE_NUM_WORKERS
,
1971 OMP_CLAUSE_VECTOR_LENGTH
};
1973 tree dims
[GOMP_DIM_MAX
];
1975 tree attr
= NULL_TREE
;
1976 unsigned non_const
= 0;
1978 for (ix
= GOMP_DIM_MAX
; ix
--;)
1980 tree clause
= omp_find_clause (clauses
, ids
[ix
]);
1981 tree dim
= NULL_TREE
;
1984 dim
= OMP_CLAUSE_EXPR (clause
, ids
[ix
]);
1986 if (dim
&& TREE_CODE (dim
) != INTEGER_CST
)
1988 dim
= integer_zero_node
;
1989 non_const
|= GOMP_DIM_MASK (ix
);
1991 attr
= tree_cons (NULL_TREE
, dim
, attr
);
1994 oacc_replace_fn_attrib (fn
, attr
);
1998 /* Push a dynamic argument set. */
1999 args
->safe_push (oacc_launch_pack (GOMP_LAUNCH_DIM
,
2000 NULL_TREE
, non_const
));
2001 for (unsigned ix
= 0; ix
!= GOMP_DIM_MAX
; ix
++)
2002 if (non_const
& GOMP_DIM_MASK (ix
))
2003 args
->safe_push (dims
[ix
]);
2007 /* Verify OpenACC routine clauses.
2009 Returns 0 if FNDECL should be marked with an OpenACC 'routine' directive, 1
2010 if it has already been marked in compatible way, and -1 if incompatible.
2011 Upon returning, the chain of clauses will contain exactly one clause
2012 specifying the level of parallelism. */
2015 oacc_verify_routine_clauses (tree fndecl
, tree
*clauses
, location_t loc
,
2016 const char *routine_str
)
2018 tree c_level
= NULL_TREE
;
2019 tree c_p
= NULL_TREE
;
2020 for (tree c
= *clauses
; c
; c_p
= c
, c
= OMP_CLAUSE_CHAIN (c
))
2021 switch (OMP_CLAUSE_CODE (c
))
2023 case OMP_CLAUSE_GANG
:
2024 case OMP_CLAUSE_WORKER
:
2025 case OMP_CLAUSE_VECTOR
:
2026 case OMP_CLAUSE_SEQ
:
2027 if (c_level
== NULL_TREE
)
2029 else if (OMP_CLAUSE_CODE (c
) == OMP_CLAUSE_CODE (c_level
))
2031 /* This has already been diagnosed in the front ends. */
2032 /* Drop the duplicate clause. */
2033 gcc_checking_assert (c_p
!= NULL_TREE
);
2034 OMP_CLAUSE_CHAIN (c_p
) = OMP_CLAUSE_CHAIN (c
);
2039 error_at (OMP_CLAUSE_LOCATION (c
),
2040 "%qs specifies a conflicting level of parallelism",
2041 omp_clause_code_name
[OMP_CLAUSE_CODE (c
)]);
2042 inform (OMP_CLAUSE_LOCATION (c_level
),
2043 "... to the previous %qs clause here",
2044 omp_clause_code_name
[OMP_CLAUSE_CODE (c_level
)]);
2045 /* Drop the conflicting clause. */
2046 gcc_checking_assert (c_p
!= NULL_TREE
);
2047 OMP_CLAUSE_CHAIN (c_p
) = OMP_CLAUSE_CHAIN (c
);
2054 if (c_level
== NULL_TREE
)
2056 /* Default to an implicit 'seq' clause. */
2057 c_level
= build_omp_clause (loc
, OMP_CLAUSE_SEQ
);
2058 OMP_CLAUSE_CHAIN (c_level
) = *clauses
;
2061 /* In *clauses, we now have exactly one clause specifying the level of
2065 = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (fndecl
));
2066 if (attr
!= NULL_TREE
)
2068 /* Diagnose if "#pragma omp declare target" has also been applied. */
2069 if (TREE_VALUE (attr
) == NULL_TREE
)
2071 /* See <https://gcc.gnu.org/PR93465>; the semantics of combining
2072 OpenACC and OpenMP 'target' are not clear. */
2074 "cannot apply %<%s%> to %qD, which has also been"
2075 " marked with an OpenMP 'declare target' directive",
2076 routine_str
, fndecl
);
2081 /* If a "#pragma acc routine" has already been applied, just verify
2082 this one for compatibility. */
2083 /* Collect previous directive's clauses. */
2084 tree c_level_p
= NULL_TREE
;
2085 for (tree c
= TREE_VALUE (attr
); c
; c
= OMP_CLAUSE_CHAIN (c
))
2086 switch (OMP_CLAUSE_CODE (c
))
2088 case OMP_CLAUSE_GANG
:
2089 case OMP_CLAUSE_WORKER
:
2090 case OMP_CLAUSE_VECTOR
:
2091 case OMP_CLAUSE_SEQ
:
2092 gcc_checking_assert (c_level_p
== NULL_TREE
);
2098 gcc_checking_assert (c_level_p
!= NULL_TREE
);
2099 /* ..., and compare to current directive's, which we've already collected
2103 /* Matching level of parallelism? */
2104 if (OMP_CLAUSE_CODE (c_level
) != OMP_CLAUSE_CODE (c_level_p
))
2107 c_diag_p
= c_level_p
;
2114 if (c_diag
!= NULL_TREE
)
2115 error_at (OMP_CLAUSE_LOCATION (c_diag
),
2116 "incompatible %qs clause when applying"
2117 " %<%s%> to %qD, which has already been"
2118 " marked with an OpenACC 'routine' directive",
2119 omp_clause_code_name
[OMP_CLAUSE_CODE (c_diag
)],
2120 routine_str
, fndecl
);
2121 else if (c_diag_p
!= NULL_TREE
)
2123 "missing %qs clause when applying"
2124 " %<%s%> to %qD, which has already been"
2125 " marked with an OpenACC 'routine' directive",
2126 omp_clause_code_name
[OMP_CLAUSE_CODE (c_diag_p
)],
2127 routine_str
, fndecl
);
2130 if (c_diag_p
!= NULL_TREE
)
2131 inform (OMP_CLAUSE_LOCATION (c_diag_p
),
2132 "... with %qs clause here",
2133 omp_clause_code_name
[OMP_CLAUSE_CODE (c_diag_p
)]);
2136 /* In the front ends, we don't preserve location information for the
2137 OpenACC routine directive itself. However, that of c_level_p
2139 location_t loc_routine
= OMP_CLAUSE_LOCATION (c_level_p
);
2140 inform (loc_routine
, "... without %qs clause near to here",
2141 omp_clause_code_name
[OMP_CLAUSE_CODE (c_diag
)]);
2150 /* Process the OpenACC 'routine' directive clauses to generate an attribute
2151 for the level of parallelism. All dimensions have a size of zero
2152 (dynamic). TREE_PURPOSE is set to indicate whether that dimension
2153 can have a loop partitioned on it. non-zero indicates
2154 yes, zero indicates no. By construction once a non-zero has been
2155 reached, further inner dimensions must also be non-zero. We set
2156 TREE_VALUE to zero for the dimensions that may be partitioned and
2157 1 for the other ones -- if a loop is (erroneously) spawned at
2158 an outer level, we don't want to try and partition it. */
2161 oacc_build_routine_dims (tree clauses
)
2163 /* Must match GOMP_DIM ordering. */
2164 static const omp_clause_code ids
[]
2165 = {OMP_CLAUSE_GANG
, OMP_CLAUSE_WORKER
, OMP_CLAUSE_VECTOR
, OMP_CLAUSE_SEQ
};
2169 for (; clauses
; clauses
= OMP_CLAUSE_CHAIN (clauses
))
2170 for (ix
= GOMP_DIM_MAX
+ 1; ix
--;)
2171 if (OMP_CLAUSE_CODE (clauses
) == ids
[ix
])
2176 gcc_checking_assert (level
>= 0);
2178 tree dims
= NULL_TREE
;
2180 for (ix
= GOMP_DIM_MAX
; ix
--;)
2181 dims
= tree_cons (build_int_cst (boolean_type_node
, ix
>= level
),
2182 build_int_cst (integer_type_node
, ix
< level
), dims
);
2187 /* Retrieve the oacc function attrib and return it. Non-oacc
2188 functions will return NULL. */
2191 oacc_get_fn_attrib (tree fn
)
2193 return lookup_attribute (OACC_FN_ATTRIB
, DECL_ATTRIBUTES (fn
));
2196 /* Return true if FN is an OpenMP or OpenACC offloading function. */
2199 offloading_function_p (tree fn
)
2201 tree attrs
= DECL_ATTRIBUTES (fn
);
2202 return (lookup_attribute ("omp declare target", attrs
)
2203 || lookup_attribute ("omp target entrypoint", attrs
));
2206 /* Extract an oacc execution dimension from FN. FN must be an
2207 offloaded function or routine that has already had its execution
2208 dimensions lowered to the target-specific values. */
2211 oacc_get_fn_dim_size (tree fn
, int axis
)
2213 tree attrs
= oacc_get_fn_attrib (fn
);
2215 gcc_assert (axis
< GOMP_DIM_MAX
);
2217 tree dims
= TREE_VALUE (attrs
);
2219 dims
= TREE_CHAIN (dims
);
2221 int size
= TREE_INT_CST_LOW (TREE_VALUE (dims
));
2226 /* Extract the dimension axis from an IFN_GOACC_DIM_POS or
2227 IFN_GOACC_DIM_SIZE call. */
2230 oacc_get_ifn_dim_arg (const gimple
*stmt
)
2232 gcc_checking_assert (gimple_call_internal_fn (stmt
) == IFN_GOACC_DIM_SIZE
2233 || gimple_call_internal_fn (stmt
) == IFN_GOACC_DIM_POS
);
2234 tree arg
= gimple_call_arg (stmt
, 0);
2235 HOST_WIDE_INT axis
= TREE_INT_CST_LOW (arg
);
2237 gcc_checking_assert (axis
>= 0 && axis
< GOMP_DIM_MAX
);
2241 #include "gt-omp-general.h"