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 (current_function_decl
646 && lookup_attribute ("omp declare target",
647 DECL_ATTRIBUTES (current_function_decl
)))
649 if (cfun
&& (cfun
->curr_properties
& PROP_gimple_any
) == 0)
651 enum tree_code construct
= OMP_TARGET
;
652 if (omp_construct_selector_matches (&construct
, 1, NULL
))
658 /* Return a name from PROP, a property in selectors accepting
662 omp_context_name_list_prop (tree prop
)
664 if (TREE_PURPOSE (prop
))
665 return IDENTIFIER_POINTER (TREE_PURPOSE (prop
));
668 const char *ret
= TREE_STRING_POINTER (TREE_VALUE (prop
));
669 if ((size_t) TREE_STRING_LENGTH (TREE_VALUE (prop
)) == strlen (ret
) + 1)
675 /* Return 1 if context selector matches the current OpenMP context, 0
676 if it does not and -1 if it is unknown and need to be determined later.
677 Some properties can be checked right away during parsing (this routine),
678 others need to wait until the whole TU is parsed, others need to wait until
679 IPA, others until vectorization. */
682 omp_context_selector_matches (tree ctx
)
685 for (tree t1
= ctx
; t1
; t1
= TREE_CHAIN (t1
))
687 char set
= IDENTIFIER_POINTER (TREE_PURPOSE (t1
))[0];
690 /* For now, ignore the construct set. While something can be
691 determined already during parsing, we don't know until end of TU
692 whether additional constructs aren't added through declare variant
693 unless "omp declare variant variant" attribute exists already
694 (so in most of the cases), and we'd need to maintain set of
695 surrounding OpenMP constructs, which is better handled during
697 if (symtab
->state
== PARSING
698 || (cfun
->curr_properties
& PROP_gimple_any
) != 0)
704 enum tree_code constructs
[5];
706 = omp_constructor_traits_to_codes (TREE_VALUE (t1
), constructs
);
707 int r
= omp_construct_selector_matches (constructs
, nconstructs
,
715 for (tree t2
= TREE_VALUE (t1
); t2
; t2
= TREE_CHAIN (t2
))
717 const char *sel
= IDENTIFIER_POINTER (TREE_PURPOSE (t2
));
721 if (set
== 'i' && !strcmp (sel
, "vendor"))
722 for (tree t3
= TREE_VALUE (t2
); t3
; t3
= TREE_CHAIN (t3
))
724 const char *prop
= omp_context_name_list_prop (t3
);
727 if ((!strcmp (prop
, " score") && TREE_PURPOSE (t3
))
728 || !strcmp (prop
, "gnu"))
734 if (set
== 'i' && !strcmp (sel
, "extension"))
735 /* We don't support any extensions right now. */
739 if (set
== 'i' && !strcmp (sel
, "atomic_default_mem_order"))
741 enum omp_memory_order omo
742 = ((enum omp_memory_order
)
744 & OMP_REQUIRES_ATOMIC_DEFAULT_MEM_ORDER
));
745 if (omo
== OMP_MEMORY_ORDER_UNSPECIFIED
)
747 /* We don't know yet, until end of TU. */
748 if (symtab
->state
== PARSING
)
754 omo
= OMP_MEMORY_ORDER_RELAXED
;
756 tree t3
= TREE_VALUE (t2
);
757 const char *prop
= IDENTIFIER_POINTER (TREE_PURPOSE (t3
));
758 if (!strcmp (prop
, " score"))
760 t3
= TREE_CHAIN (t3
);
761 prop
= IDENTIFIER_POINTER (TREE_PURPOSE (t3
));
763 if (!strcmp (prop
, "relaxed")
764 && omo
!= OMP_MEMORY_ORDER_RELAXED
)
766 else if (!strcmp (prop
, "seq_cst")
767 && omo
!= OMP_MEMORY_ORDER_SEQ_CST
)
769 else if (!strcmp (prop
, "acq_rel")
770 && omo
!= OMP_MEMORY_ORDER_ACQ_REL
)
773 if (set
== 'd' && !strcmp (sel
, "arch"))
774 for (tree t3
= TREE_VALUE (t2
); t3
; t3
= TREE_CHAIN (t3
))
776 const char *arch
= omp_context_name_list_prop (t3
);
780 if (targetm
.omp
.device_kind_arch_isa
!= NULL
)
781 r
= targetm
.omp
.device_kind_arch_isa (omp_device_arch
,
783 if (r
== 0 || (r
== -1 && symtab
->state
!= PARSING
))
785 /* If we are or might be in a target region or
786 declare target function, need to take into account
787 also offloading values. */
788 if (!omp_maybe_offloaded ())
790 if (strcmp (arch
, "hsa") == 0
791 && hsa_gen_requested_p ())
796 if (ENABLE_OFFLOADING
)
798 const char *arches
= omp_offload_device_arch
;
799 if (omp_offload_device_kind_arch_isa (arches
,
810 /* If arch matches on the host, it still might not match
811 in the offloading region. */
812 else if (omp_maybe_offloaded ())
817 if (set
== 'i' && !strcmp (sel
, "unified_address"))
819 if ((omp_requires_mask
& OMP_REQUIRES_UNIFIED_ADDRESS
) == 0)
821 if (symtab
->state
== PARSING
)
828 if (set
== 'i' && !strcmp (sel
, "unified_shared_memory"))
830 if ((omp_requires_mask
831 & OMP_REQUIRES_UNIFIED_SHARED_MEMORY
) == 0)
833 if (symtab
->state
== PARSING
)
842 if (set
== 'i' && !strcmp (sel
, "dynamic_allocators"))
844 if ((omp_requires_mask
845 & OMP_REQUIRES_DYNAMIC_ALLOCATORS
) == 0)
847 if (symtab
->state
== PARSING
)
856 if (set
== 'i' && !strcmp (sel
, "reverse_offload"))
858 if ((omp_requires_mask
& OMP_REQUIRES_REVERSE_OFFLOAD
) == 0)
860 if (symtab
->state
== PARSING
)
869 if (set
== 'd' && !strcmp (sel
, "kind"))
870 for (tree t3
= TREE_VALUE (t2
); t3
; t3
= TREE_CHAIN (t3
))
872 const char *prop
= omp_context_name_list_prop (t3
);
875 if (!strcmp (prop
, "any"))
877 if (!strcmp (prop
, "host"))
879 if (omp_maybe_offloaded ())
883 if (!strcmp (prop
, "nohost"))
885 if (omp_maybe_offloaded ())
892 if (targetm
.omp
.device_kind_arch_isa
!= NULL
)
893 r
= targetm
.omp
.device_kind_arch_isa (omp_device_kind
,
896 r
= strcmp (prop
, "cpu") == 0;
897 if (r
== 0 || (r
== -1 && symtab
->state
!= PARSING
))
899 /* If we are or might be in a target region or
900 declare target function, need to take into account
901 also offloading values. */
902 if (!omp_maybe_offloaded ())
904 if (strcmp (prop
, "gpu") == 0
905 && hsa_gen_requested_p ())
910 if (ENABLE_OFFLOADING
)
912 const char *kinds
= omp_offload_device_kind
;
913 if (omp_offload_device_kind_arch_isa (kinds
, prop
))
923 /* If kind matches on the host, it still might not match
924 in the offloading region. */
925 else if (omp_maybe_offloaded ())
930 if (set
== 'd' && !strcmp (sel
, "isa"))
931 for (tree t3
= TREE_VALUE (t2
); t3
; t3
= TREE_CHAIN (t3
))
933 const char *isa
= omp_context_name_list_prop (t3
);
937 if (targetm
.omp
.device_kind_arch_isa
!= NULL
)
938 r
= targetm
.omp
.device_kind_arch_isa (omp_device_isa
,
940 if (r
== 0 || (r
== -1 && symtab
->state
!= PARSING
))
942 /* If isa is valid on the target, but not in the
943 current function and current function has
944 #pragma omp declare simd on it, some simd clones
945 might have the isa added later on. */
947 && targetm
.simd_clone
.compute_vecsize_and_simdlen
)
950 = DECL_ATTRIBUTES (current_function_decl
);
951 if (lookup_attribute ("omp declare simd", attrs
))
957 /* If we are or might be in a target region or
958 declare target function, need to take into account
959 also offloading values. */
960 if (!omp_maybe_offloaded ())
962 if (ENABLE_OFFLOADING
)
964 const char *isas
= omp_offload_device_isa
;
965 if (omp_offload_device_kind_arch_isa (isas
, isa
))
975 /* If isa matches on the host, it still might not match
976 in the offloading region. */
977 else if (omp_maybe_offloaded ())
982 if (set
== 'u' && !strcmp (sel
, "condition"))
983 for (tree t3
= TREE_VALUE (t2
); t3
; t3
= TREE_CHAIN (t3
))
984 if (TREE_PURPOSE (t3
) == NULL_TREE
)
986 if (integer_zerop (TREE_VALUE (t3
)))
988 if (integer_nonzerop (TREE_VALUE (t3
)))
1001 /* Compare construct={simd} CLAUSES1 with CLAUSES2, return 0/-1/1/2 as
1002 in omp_context_selector_set_compare. */
1005 omp_construct_simd_compare (tree clauses1
, tree clauses2
)
1007 if (clauses1
== NULL_TREE
)
1008 return clauses2
== NULL_TREE
? 0 : -1;
1009 if (clauses2
== NULL_TREE
)
1013 struct declare_variant_simd_data
{
1014 bool inbranch
, notinbranch
;
1016 auto_vec
<tree
,16> data_sharing
;
1017 auto_vec
<tree
,16> aligned
;
1018 declare_variant_simd_data ()
1019 : inbranch(false), notinbranch(false), simdlen(NULL_TREE
) {}
1022 for (i
= 0; i
< 2; i
++)
1023 for (tree c
= i
? clauses2
: clauses1
; c
; c
= OMP_CLAUSE_CHAIN (c
))
1026 switch (OMP_CLAUSE_CODE (c
))
1028 case OMP_CLAUSE_INBRANCH
:
1029 data
[i
].inbranch
= true;
1031 case OMP_CLAUSE_NOTINBRANCH
:
1032 data
[i
].notinbranch
= true;
1034 case OMP_CLAUSE_SIMDLEN
:
1035 data
[i
].simdlen
= OMP_CLAUSE_SIMDLEN_EXPR (c
);
1037 case OMP_CLAUSE_UNIFORM
:
1038 case OMP_CLAUSE_LINEAR
:
1039 v
= &data
[i
].data_sharing
;
1041 case OMP_CLAUSE_ALIGNED
:
1042 v
= &data
[i
].aligned
;
1047 unsigned HOST_WIDE_INT argno
= tree_to_uhwi (OMP_CLAUSE_DECL (c
));
1048 if (argno
>= v
->length ())
1049 v
->safe_grow_cleared (argno
+ 1);
1052 /* Here, r is used as a bitmask, 2 is set if CLAUSES1 has something
1053 CLAUSES2 doesn't, 1 is set if CLAUSES2 has something CLAUSES1
1054 doesn't. Thus, r == 3 implies return value 2, r == 1 implies
1055 -1, r == 2 implies 1 and r == 0 implies 0. */
1056 if (data
[0].inbranch
!= data
[1].inbranch
)
1057 r
|= data
[0].inbranch
? 2 : 1;
1058 if (data
[0].notinbranch
!= data
[1].notinbranch
)
1059 r
|= data
[0].notinbranch
? 2 : 1;
1060 if (!simple_cst_equal (data
[0].simdlen
, data
[1].simdlen
))
1062 if (data
[0].simdlen
&& data
[1].simdlen
)
1064 r
|= data
[0].simdlen
? 2 : 1;
1066 if (data
[0].data_sharing
.length () < data
[1].data_sharing
.length ()
1067 || data
[0].aligned
.length () < data
[1].aligned
.length ())
1070 FOR_EACH_VEC_ELT (data
[0].data_sharing
, i
, c1
)
1072 c2
= (i
< data
[1].data_sharing
.length ()
1073 ? data
[1].data_sharing
[i
] : NULL_TREE
);
1074 if ((c1
== NULL_TREE
) != (c2
== NULL_TREE
))
1076 r
|= c1
!= NULL_TREE
? 2 : 1;
1079 if (c1
== NULL_TREE
)
1081 if (OMP_CLAUSE_CODE (c1
) != OMP_CLAUSE_CODE (c2
))
1083 if (OMP_CLAUSE_CODE (c1
) != OMP_CLAUSE_LINEAR
)
1085 if (OMP_CLAUSE_LINEAR_VARIABLE_STRIDE (c1
)
1086 != OMP_CLAUSE_LINEAR_VARIABLE_STRIDE (c2
))
1088 if (OMP_CLAUSE_LINEAR_KIND (c1
) != OMP_CLAUSE_LINEAR_KIND (c2
))
1090 if (!simple_cst_equal (OMP_CLAUSE_LINEAR_STEP (c1
),
1091 OMP_CLAUSE_LINEAR_STEP (c2
)))
1094 FOR_EACH_VEC_ELT (data
[0].aligned
, i
, c1
)
1096 c2
= i
< data
[1].aligned
.length () ? data
[1].aligned
[i
] : NULL_TREE
;
1097 if ((c1
== NULL_TREE
) != (c2
== NULL_TREE
))
1099 r
|= c1
!= NULL_TREE
? 2 : 1;
1102 if (c1
== NULL_TREE
)
1104 if (!simple_cst_equal (OMP_CLAUSE_ALIGNED_ALIGNMENT (c1
),
1105 OMP_CLAUSE_ALIGNED_ALIGNMENT (c2
)))
1114 default: gcc_unreachable ();
1118 /* Compare properties of selectors SEL from SET other than construct.
1119 Return 0/-1/1/2 as in omp_context_selector_set_compare.
1120 Unlike set names or selector names, properties can have duplicates. */
1123 omp_context_selector_props_compare (const char *set
, const char *sel
,
1124 tree ctx1
, tree ctx2
)
1127 for (int pass
= 0; pass
< 2; pass
++)
1128 for (tree t1
= pass
? ctx2
: ctx1
; t1
; t1
= TREE_CHAIN (t1
))
1131 for (t2
= pass
? ctx1
: ctx2
; t2
; t2
= TREE_CHAIN (t2
))
1132 if (TREE_PURPOSE (t1
) == TREE_PURPOSE (t2
))
1134 if (TREE_PURPOSE (t1
) == NULL_TREE
)
1136 if (set
[0] == 'u' && strcmp (sel
, "condition") == 0)
1138 if (integer_zerop (TREE_VALUE (t1
))
1139 != integer_zerop (TREE_VALUE (t2
)))
1143 if (simple_cst_equal (TREE_VALUE (t1
), TREE_VALUE (t2
)))
1146 else if (strcmp (IDENTIFIER_POINTER (TREE_PURPOSE (t1
)),
1149 if (!simple_cst_equal (TREE_VALUE (t1
), TREE_VALUE (t2
)))
1156 else if (TREE_PURPOSE (t1
)
1157 && TREE_PURPOSE (t2
) == NULL_TREE
1158 && TREE_CODE (TREE_VALUE (t2
)) == STRING_CST
)
1160 const char *p1
= omp_context_name_list_prop (t1
);
1161 const char *p2
= omp_context_name_list_prop (t2
);
1163 && strcmp (p1
, p2
) == 0
1164 && strcmp (p1
, " score"))
1167 else if (TREE_PURPOSE (t1
) == NULL_TREE
1168 && TREE_PURPOSE (t2
)
1169 && TREE_CODE (TREE_VALUE (t1
)) == STRING_CST
)
1171 const char *p1
= omp_context_name_list_prop (t1
);
1172 const char *p2
= omp_context_name_list_prop (t2
);
1174 && strcmp (p1
, p2
) == 0
1175 && strcmp (p1
, " score"))
1178 if (t2
== NULL_TREE
)
1180 int r
= pass
? -1 : 1;
1181 if (ret
&& ret
!= r
)
1195 /* Compare single context selector sets CTX1 and CTX2 with SET name.
1196 Return 0 if CTX1 is equal to CTX2,
1197 -1 if CTX1 is a strict subset of CTX2,
1198 1 if CTX2 is a strict subset of CTX1, or
1199 2 if neither context is a subset of another one. */
1202 omp_context_selector_set_compare (const char *set
, tree ctx1
, tree ctx2
)
1204 bool swapped
= false;
1206 int len1
= list_length (ctx1
);
1207 int len2
= list_length (ctx2
);
1212 std::swap (ctx1
, ctx2
);
1213 std::swap (len1
, len2
);
1219 tree simd
= get_identifier ("simd");
1220 /* Handle construct set specially. In this case the order
1221 of the selector matters too. */
1222 for (t1
= ctx1
; t1
; t1
= TREE_CHAIN (t1
))
1223 if (TREE_PURPOSE (t1
) == TREE_PURPOSE (t2
))
1226 if (TREE_PURPOSE (t1
) == simd
)
1227 r
= omp_construct_simd_compare (TREE_VALUE (t1
),
1229 if (r
== 2 || (ret
&& r
&& (ret
< 0) != (r
< 0)))
1233 t2
= TREE_CHAIN (t2
);
1234 if (t2
== NULL_TREE
)
1236 t1
= TREE_CHAIN (t1
);
1244 if (t2
!= NULL_TREE
)
1246 if (t1
!= NULL_TREE
)
1254 return swapped
? -ret
: ret
;
1256 for (tree t1
= ctx1
; t1
; t1
= TREE_CHAIN (t1
))
1259 for (t2
= ctx2
; t2
; t2
= TREE_CHAIN (t2
))
1260 if (TREE_PURPOSE (t1
) == TREE_PURPOSE (t2
))
1262 const char *sel
= IDENTIFIER_POINTER (TREE_PURPOSE (t1
));
1263 int r
= omp_context_selector_props_compare (set
, sel
,
1266 if (r
== 2 || (ret
&& r
&& (ret
< 0) != (r
< 0)))
1273 if (t2
== NULL_TREE
)
1284 return swapped
? -ret
: ret
;
1287 /* Compare whole context selector specification CTX1 and CTX2.
1288 Return 0 if CTX1 is equal to CTX2,
1289 -1 if CTX1 is a strict subset of CTX2,
1290 1 if CTX2 is a strict subset of CTX1, or
1291 2 if neither context is a subset of another one. */
1294 omp_context_selector_compare (tree ctx1
, tree ctx2
)
1296 bool swapped
= false;
1298 int len1
= list_length (ctx1
);
1299 int len2
= list_length (ctx2
);
1304 std::swap (ctx1
, ctx2
);
1305 std::swap (len1
, len2
);
1307 for (tree t1
= ctx1
; t1
; t1
= TREE_CHAIN (t1
))
1310 for (t2
= ctx2
; t2
; t2
= TREE_CHAIN (t2
))
1311 if (TREE_PURPOSE (t1
) == TREE_PURPOSE (t2
))
1313 const char *set
= IDENTIFIER_POINTER (TREE_PURPOSE (t1
));
1314 int r
= omp_context_selector_set_compare (set
, TREE_VALUE (t1
),
1316 if (r
== 2 || (ret
&& r
&& (ret
< 0) != (r
< 0)))
1323 if (t2
== NULL_TREE
)
1334 return swapped
? -ret
: ret
;
1337 /* From context selector CTX, return trait-selector with name SEL in
1338 trait-selector-set with name SET if any, or NULL_TREE if not found.
1339 If SEL is NULL, return the list of trait-selectors in SET. */
1342 omp_get_context_selector (tree ctx
, const char *set
, const char *sel
)
1344 tree setid
= get_identifier (set
);
1345 tree selid
= sel
? get_identifier (sel
) : NULL_TREE
;
1346 for (tree t1
= ctx
; t1
; t1
= TREE_CHAIN (t1
))
1347 if (TREE_PURPOSE (t1
) == setid
)
1350 return TREE_VALUE (t1
);
1351 for (tree t2
= TREE_VALUE (t1
); t2
; t2
= TREE_CHAIN (t2
))
1352 if (TREE_PURPOSE (t2
) == selid
)
1358 /* Compute *SCORE for context selector CTX. Return true if the score
1359 would be different depending on whether it is a declare simd clone or
1360 not. DECLARE_SIMD should be true for the case when it would be
1361 a declare simd clone. */
1364 omp_context_compute_score (tree ctx
, widest_int
*score
, bool declare_simd
)
1366 tree construct
= omp_get_context_selector (ctx
, "construct", NULL
);
1367 bool has_kind
= omp_get_context_selector (ctx
, "device", "kind");
1368 bool has_arch
= omp_get_context_selector (ctx
, "device", "arch");
1369 bool has_isa
= omp_get_context_selector (ctx
, "device", "isa");
1372 for (tree t1
= ctx
; t1
; t1
= TREE_CHAIN (t1
))
1373 if (TREE_VALUE (t1
) != construct
)
1374 for (tree t2
= TREE_VALUE (t1
); t2
; t2
= TREE_CHAIN (t2
))
1375 if (tree t3
= TREE_VALUE (t2
))
1376 if (TREE_PURPOSE (t3
)
1377 && strcmp (IDENTIFIER_POINTER (TREE_PURPOSE (t3
)), " score") == 0
1378 && TREE_CODE (TREE_VALUE (t3
)) == INTEGER_CST
)
1379 *score
+= wi::to_widest (TREE_VALUE (t3
));
1380 if (construct
|| has_kind
|| has_arch
|| has_isa
)
1383 enum tree_code constructs
[5];
1384 int nconstructs
= 0;
1386 nconstructs
= omp_constructor_traits_to_codes (construct
, constructs
);
1387 if (omp_construct_selector_matches (constructs
, nconstructs
, scores
)
1390 int b
= declare_simd
? nconstructs
+ 1 : 0;
1391 if (scores
[b
+ nconstructs
] + 4U < score
->get_precision ())
1393 for (int n
= 0; n
< nconstructs
; ++n
)
1395 if (scores
[b
+ n
] < 0)
1400 *score
+= wi::shifted_mask
<widest_int
> (scores
[b
+ n
], 1, false);
1403 *score
+= wi::shifted_mask
<widest_int
> (scores
[b
+ nconstructs
],
1406 *score
+= wi::shifted_mask
<widest_int
> (scores
[b
+ nconstructs
] + 1,
1409 *score
+= wi::shifted_mask
<widest_int
> (scores
[b
+ nconstructs
] + 2,
1412 else /* FIXME: Implement this. */
1418 /* Try to resolve declare variant, return the variant decl if it should
1419 be used instead of base, or base otherwise. */
1422 omp_resolve_declare_variant (tree base
)
1424 tree variant1
= NULL_TREE
, variant2
= NULL_TREE
;
1425 auto_vec
<tree
, 16> variants
;
1426 auto_vec
<bool, 16> defer
;
1427 bool any_deferred
= false;
1428 for (tree attr
= DECL_ATTRIBUTES (base
); attr
; attr
= TREE_CHAIN (attr
))
1430 attr
= lookup_attribute ("omp declare variant base", attr
);
1431 if (attr
== NULL_TREE
)
1433 if (TREE_CODE (TREE_PURPOSE (TREE_VALUE (attr
))) != FUNCTION_DECL
)
1435 switch (omp_context_selector_matches (TREE_VALUE (TREE_VALUE (attr
))))
1438 /* No match, ignore. */
1441 /* Needs to be deferred. */
1442 any_deferred
= true;
1443 variants
.safe_push (attr
);
1444 defer
.safe_push (true);
1447 variants
.safe_push (attr
);
1448 defer
.safe_push (false);
1452 if (variants
.length () == 0)
1457 widest_int max_score1
= 0;
1458 widest_int max_score2
= 0;
1462 FOR_EACH_VEC_ELT (variants
, i
, attr1
)
1467 tree ctx
= TREE_VALUE (TREE_VALUE (attr1
));
1468 need_two
= omp_context_compute_score (ctx
, &score1
, false);
1470 omp_context_compute_score (ctx
, &score2
, true);
1476 max_score1
= score1
;
1477 max_score2
= score2
;
1486 if (max_score1
== score1
)
1487 variant1
= NULL_TREE
;
1488 else if (score1
> max_score1
)
1490 max_score1
= score1
;
1491 variant1
= defer
[i
] ? NULL_TREE
: attr1
;
1493 if (max_score2
== score2
)
1494 variant2
= NULL_TREE
;
1495 else if (score2
> max_score2
)
1497 max_score2
= score2
;
1498 variant2
= defer
[i
] ? NULL_TREE
: attr1
;
1503 /* If there is a clear winner variant with the score which is not
1504 deferred, verify it is not a strict subset of any other context
1505 selector and if it is not, it is the best alternative no matter
1506 whether the others do or don't match. */
1507 if (variant1
&& variant1
== variant2
)
1509 tree ctx1
= TREE_VALUE (TREE_VALUE (variant1
));
1510 FOR_EACH_VEC_ELT (variants
, i
, attr2
)
1512 if (attr2
== variant1
)
1514 tree ctx2
= TREE_VALUE (TREE_VALUE (attr2
));
1515 int r
= omp_context_selector_compare (ctx1
, ctx2
);
1518 /* The winner is a strict subset of ctx2, can't
1520 variant1
= NULL_TREE
;
1525 return TREE_PURPOSE (TREE_VALUE (variant1
));
1531 if (variants
.length () == 1)
1532 return TREE_PURPOSE (TREE_VALUE (variants
[0]));
1534 /* A context selector that is a strict subset of another context selector has a score
1538 FOR_EACH_VEC_ELT (variants
, i
, attr1
)
1541 tree ctx1
= TREE_VALUE (TREE_VALUE (attr1
));
1542 FOR_EACH_VEC_ELT_FROM (variants
, j
, attr2
, i
+ 1)
1545 tree ctx2
= TREE_VALUE (TREE_VALUE (attr2
));
1546 int r
= omp_context_selector_compare (ctx1
, ctx2
);
1549 /* ctx1 is a strict subset of ctx2, remove
1550 attr1 from the vector. */
1551 variants
[i
] = NULL_TREE
;
1555 /* ctx2 is a strict subset of ctx1, remove attr2
1557 variants
[j
] = NULL_TREE
;
1560 widest_int max_score1
= 0;
1561 widest_int max_score2
= 0;
1563 FOR_EACH_VEC_ELT (variants
, i
, attr1
)
1575 ctx
= TREE_VALUE (TREE_VALUE (variant1
));
1576 need_two
= omp_context_compute_score (ctx
, &max_score1
, false);
1578 omp_context_compute_score (ctx
, &max_score2
, true);
1580 max_score2
= max_score1
;
1582 ctx
= TREE_VALUE (TREE_VALUE (attr1
));
1583 need_two
= omp_context_compute_score (ctx
, &score1
, false);
1585 omp_context_compute_score (ctx
, &score2
, true);
1588 if (score1
> max_score1
)
1590 max_score1
= score1
;
1593 if (score2
> max_score2
)
1595 max_score2
= score2
;
1605 /* If there is a disagreement on which variant has the highest score
1606 depending on whether it will be in a declare simd clone or not,
1607 punt for now and defer until after IPA where we will know that. */
1608 return ((variant1
&& variant1
== variant2
)
1609 ? TREE_PURPOSE (TREE_VALUE (variant1
)) : base
);
1613 /* Encode an oacc launch argument. This matches the GOMP_LAUNCH_PACK
1614 macro on gomp-constants.h. We do not check for overflow. */
1617 oacc_launch_pack (unsigned code
, tree device
, unsigned op
)
1621 res
= build_int_cst (unsigned_type_node
, GOMP_LAUNCH_PACK (code
, 0, op
));
1624 device
= fold_build2 (LSHIFT_EXPR
, unsigned_type_node
,
1625 device
, build_int_cst (unsigned_type_node
,
1626 GOMP_LAUNCH_DEVICE_SHIFT
));
1627 res
= fold_build2 (BIT_IOR_EXPR
, unsigned_type_node
, res
, device
);
1632 /* FIXME: What is the following comment for? */
1633 /* Look for compute grid dimension clauses and convert to an attribute
1634 attached to FN. This permits the target-side code to (a) massage
1635 the dimensions, (b) emit that data and (c) optimize. Non-constant
1636 dimensions are pushed onto ARGS.
1638 The attribute value is a TREE_LIST. A set of dimensions is
1639 represented as a list of INTEGER_CST. Those that are runtime
1640 exprs are represented as an INTEGER_CST of zero.
1642 TODO: Normally the attribute will just contain a single such list. If
1643 however it contains a list of lists, this will represent the use of
1644 device_type. Each member of the outer list is an assoc list of
1645 dimensions, keyed by the device type. The first entry will be the
1646 default. Well, that's the plan. */
1648 /* Replace any existing oacc fn attribute with updated dimensions. */
1650 /* Variant working on a list of attributes. */
1653 oacc_replace_fn_attrib_attr (tree attribs
, tree dims
)
1655 tree ident
= get_identifier (OACC_FN_ATTRIB
);
1657 /* If we happen to be present as the first attrib, drop it. */
1658 if (attribs
&& TREE_PURPOSE (attribs
) == ident
)
1659 attribs
= TREE_CHAIN (attribs
);
1660 return tree_cons (ident
, dims
, attribs
);
1663 /* Variant working on a function decl. */
1666 oacc_replace_fn_attrib (tree fn
, tree dims
)
1668 DECL_ATTRIBUTES (fn
)
1669 = oacc_replace_fn_attrib_attr (DECL_ATTRIBUTES (fn
), dims
);
1672 /* Scan CLAUSES for launch dimensions and attach them to the oacc
1673 function attribute. Push any that are non-constant onto the ARGS
1674 list, along with an appropriate GOMP_LAUNCH_DIM tag. */
1677 oacc_set_fn_attrib (tree fn
, tree clauses
, vec
<tree
> *args
)
1679 /* Must match GOMP_DIM ordering. */
1680 static const omp_clause_code ids
[]
1681 = { OMP_CLAUSE_NUM_GANGS
, OMP_CLAUSE_NUM_WORKERS
,
1682 OMP_CLAUSE_VECTOR_LENGTH
};
1684 tree dims
[GOMP_DIM_MAX
];
1686 tree attr
= NULL_TREE
;
1687 unsigned non_const
= 0;
1689 for (ix
= GOMP_DIM_MAX
; ix
--;)
1691 tree clause
= omp_find_clause (clauses
, ids
[ix
]);
1692 tree dim
= NULL_TREE
;
1695 dim
= OMP_CLAUSE_EXPR (clause
, ids
[ix
]);
1697 if (dim
&& TREE_CODE (dim
) != INTEGER_CST
)
1699 dim
= integer_zero_node
;
1700 non_const
|= GOMP_DIM_MASK (ix
);
1702 attr
= tree_cons (NULL_TREE
, dim
, attr
);
1705 oacc_replace_fn_attrib (fn
, attr
);
1709 /* Push a dynamic argument set. */
1710 args
->safe_push (oacc_launch_pack (GOMP_LAUNCH_DIM
,
1711 NULL_TREE
, non_const
));
1712 for (unsigned ix
= 0; ix
!= GOMP_DIM_MAX
; ix
++)
1713 if (non_const
& GOMP_DIM_MASK (ix
))
1714 args
->safe_push (dims
[ix
]);
1718 /* Verify OpenACC routine clauses.
1720 Returns 0 if FNDECL should be marked with an OpenACC 'routine' directive, 1
1721 if it has already been marked in compatible way, and -1 if incompatible.
1722 Upon returning, the chain of clauses will contain exactly one clause
1723 specifying the level of parallelism. */
1726 oacc_verify_routine_clauses (tree fndecl
, tree
*clauses
, location_t loc
,
1727 const char *routine_str
)
1729 tree c_level
= NULL_TREE
;
1730 tree c_p
= NULL_TREE
;
1731 for (tree c
= *clauses
; c
; c_p
= c
, c
= OMP_CLAUSE_CHAIN (c
))
1732 switch (OMP_CLAUSE_CODE (c
))
1734 case OMP_CLAUSE_GANG
:
1735 case OMP_CLAUSE_WORKER
:
1736 case OMP_CLAUSE_VECTOR
:
1737 case OMP_CLAUSE_SEQ
:
1738 if (c_level
== NULL_TREE
)
1740 else if (OMP_CLAUSE_CODE (c
) == OMP_CLAUSE_CODE (c_level
))
1742 /* This has already been diagnosed in the front ends. */
1743 /* Drop the duplicate clause. */
1744 gcc_checking_assert (c_p
!= NULL_TREE
);
1745 OMP_CLAUSE_CHAIN (c_p
) = OMP_CLAUSE_CHAIN (c
);
1750 error_at (OMP_CLAUSE_LOCATION (c
),
1751 "%qs specifies a conflicting level of parallelism",
1752 omp_clause_code_name
[OMP_CLAUSE_CODE (c
)]);
1753 inform (OMP_CLAUSE_LOCATION (c_level
),
1754 "... to the previous %qs clause here",
1755 omp_clause_code_name
[OMP_CLAUSE_CODE (c_level
)]);
1756 /* Drop the conflicting clause. */
1757 gcc_checking_assert (c_p
!= NULL_TREE
);
1758 OMP_CLAUSE_CHAIN (c_p
) = OMP_CLAUSE_CHAIN (c
);
1765 if (c_level
== NULL_TREE
)
1767 /* Default to an implicit 'seq' clause. */
1768 c_level
= build_omp_clause (loc
, OMP_CLAUSE_SEQ
);
1769 OMP_CLAUSE_CHAIN (c_level
) = *clauses
;
1772 /* In *clauses, we now have exactly one clause specifying the level of
1776 = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (fndecl
));
1777 if (attr
!= NULL_TREE
)
1779 /* If a "#pragma acc routine" has already been applied, just verify
1780 this one for compatibility. */
1781 /* Collect previous directive's clauses. */
1782 tree c_level_p
= NULL_TREE
;
1783 for (tree c
= TREE_VALUE (attr
); c
; c
= OMP_CLAUSE_CHAIN (c
))
1784 switch (OMP_CLAUSE_CODE (c
))
1786 case OMP_CLAUSE_GANG
:
1787 case OMP_CLAUSE_WORKER
:
1788 case OMP_CLAUSE_VECTOR
:
1789 case OMP_CLAUSE_SEQ
:
1790 gcc_checking_assert (c_level_p
== NULL_TREE
);
1796 gcc_checking_assert (c_level_p
!= NULL_TREE
);
1797 /* ..., and compare to current directive's, which we've already collected
1801 /* Matching level of parallelism? */
1802 if (OMP_CLAUSE_CODE (c_level
) != OMP_CLAUSE_CODE (c_level_p
))
1805 c_diag_p
= c_level_p
;
1812 if (c_diag
!= NULL_TREE
)
1813 error_at (OMP_CLAUSE_LOCATION (c_diag
),
1814 "incompatible %qs clause when applying"
1815 " %<%s%> to %qD, which has already been"
1816 " marked with an OpenACC 'routine' directive",
1817 omp_clause_code_name
[OMP_CLAUSE_CODE (c_diag
)],
1818 routine_str
, fndecl
);
1819 else if (c_diag_p
!= NULL_TREE
)
1821 "missing %qs clause when applying"
1822 " %<%s%> to %qD, which has already been"
1823 " marked with an OpenACC 'routine' directive",
1824 omp_clause_code_name
[OMP_CLAUSE_CODE (c_diag_p
)],
1825 routine_str
, fndecl
);
1828 if (c_diag_p
!= NULL_TREE
)
1829 inform (OMP_CLAUSE_LOCATION (c_diag_p
),
1830 "... with %qs clause here",
1831 omp_clause_code_name
[OMP_CLAUSE_CODE (c_diag_p
)]);
1834 /* In the front ends, we don't preserve location information for the
1835 OpenACC routine directive itself. However, that of c_level_p
1837 location_t loc_routine
= OMP_CLAUSE_LOCATION (c_level_p
);
1838 inform (loc_routine
, "... without %qs clause near to here",
1839 omp_clause_code_name
[OMP_CLAUSE_CODE (c_diag
)]);
1848 /* Process the OpenACC 'routine' directive clauses to generate an attribute
1849 for the level of parallelism. All dimensions have a size of zero
1850 (dynamic). TREE_PURPOSE is set to indicate whether that dimension
1851 can have a loop partitioned on it. non-zero indicates
1852 yes, zero indicates no. By construction once a non-zero has been
1853 reached, further inner dimensions must also be non-zero. We set
1854 TREE_VALUE to zero for the dimensions that may be partitioned and
1855 1 for the other ones -- if a loop is (erroneously) spawned at
1856 an outer level, we don't want to try and partition it. */
1859 oacc_build_routine_dims (tree clauses
)
1861 /* Must match GOMP_DIM ordering. */
1862 static const omp_clause_code ids
[]
1863 = {OMP_CLAUSE_GANG
, OMP_CLAUSE_WORKER
, OMP_CLAUSE_VECTOR
, OMP_CLAUSE_SEQ
};
1867 for (; clauses
; clauses
= OMP_CLAUSE_CHAIN (clauses
))
1868 for (ix
= GOMP_DIM_MAX
+ 1; ix
--;)
1869 if (OMP_CLAUSE_CODE (clauses
) == ids
[ix
])
1874 gcc_checking_assert (level
>= 0);
1876 tree dims
= NULL_TREE
;
1878 for (ix
= GOMP_DIM_MAX
; ix
--;)
1879 dims
= tree_cons (build_int_cst (boolean_type_node
, ix
>= level
),
1880 build_int_cst (integer_type_node
, ix
< level
), dims
);
1885 /* Retrieve the oacc function attrib and return it. Non-oacc
1886 functions will return NULL. */
1889 oacc_get_fn_attrib (tree fn
)
1891 return lookup_attribute (OACC_FN_ATTRIB
, DECL_ATTRIBUTES (fn
));
1894 /* Return true if FN is an OpenMP or OpenACC offloading function. */
1897 offloading_function_p (tree fn
)
1899 tree attrs
= DECL_ATTRIBUTES (fn
);
1900 return (lookup_attribute ("omp declare target", attrs
)
1901 || lookup_attribute ("omp target entrypoint", attrs
));
1904 /* Extract an oacc execution dimension from FN. FN must be an
1905 offloaded function or routine that has already had its execution
1906 dimensions lowered to the target-specific values. */
1909 oacc_get_fn_dim_size (tree fn
, int axis
)
1911 tree attrs
= oacc_get_fn_attrib (fn
);
1913 gcc_assert (axis
< GOMP_DIM_MAX
);
1915 tree dims
= TREE_VALUE (attrs
);
1917 dims
= TREE_CHAIN (dims
);
1919 int size
= TREE_INT_CST_LOW (TREE_VALUE (dims
));
1924 /* Extract the dimension axis from an IFN_GOACC_DIM_POS or
1925 IFN_GOACC_DIM_SIZE call. */
1928 oacc_get_ifn_dim_arg (const gimple
*stmt
)
1930 gcc_checking_assert (gimple_call_internal_fn (stmt
) == IFN_GOACC_DIM_SIZE
1931 || gimple_call_internal_fn (stmt
) == IFN_GOACC_DIM_POS
);
1932 tree arg
= gimple_call_arg (stmt
, 0);
1933 HOST_WIDE_INT axis
= TREE_INT_CST_LOW (arg
);
1935 gcc_checking_assert (axis
>= 0 && axis
< GOMP_DIM_MAX
);