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"
45 #include "tree-iterator.h"
47 enum omp_requires omp_requires_mask
;
50 omp_find_clause (tree clauses
, enum omp_clause_code kind
)
52 for (; clauses
; clauses
= OMP_CLAUSE_CHAIN (clauses
))
53 if (OMP_CLAUSE_CODE (clauses
) == kind
)
59 /* True if OpenMP should regard this DECL as being a scalar which has Fortran's
60 allocatable or pointer attribute. */
62 omp_is_allocatable_or_ptr (tree decl
)
64 return lang_hooks
.decls
.omp_is_allocatable_or_ptr (decl
);
67 /* Check whether this DECL belongs to a Fortran optional argument.
68 With 'for_present_check' set to false, decls which are optional parameters
69 themselve are returned as tree - or a NULL_TREE otherwise. Those decls are
70 always pointers. With 'for_present_check' set to true, the decl for checking
71 whether an argument is present is returned; for arguments with value
72 attribute this is the hidden argument and of BOOLEAN_TYPE. If the decl is
73 unrelated to optional arguments, NULL_TREE is returned. */
76 omp_check_optional_argument (tree decl
, bool for_present_check
)
78 return lang_hooks
.decls
.omp_check_optional_argument (decl
, for_present_check
);
81 /* Return true if DECL is a reference type. */
84 omp_is_reference (tree decl
)
86 return lang_hooks
.decls
.omp_privatize_by_reference (decl
);
89 /* Adjust *COND_CODE and *N2 so that the former is either LT_EXPR or GT_EXPR,
90 given that V is the loop index variable and STEP is loop step. */
93 omp_adjust_for_condition (location_t loc
, enum tree_code
*cond_code
, tree
*n2
,
103 gcc_assert (TREE_CODE (step
) == INTEGER_CST
);
104 if (TREE_CODE (TREE_TYPE (v
)) == INTEGER_TYPE
)
106 if (integer_onep (step
))
107 *cond_code
= LT_EXPR
;
110 gcc_assert (integer_minus_onep (step
));
111 *cond_code
= GT_EXPR
;
116 tree unit
= TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (v
)));
117 gcc_assert (TREE_CODE (unit
) == INTEGER_CST
);
118 if (tree_int_cst_equal (unit
, step
))
119 *cond_code
= LT_EXPR
;
122 gcc_assert (wi::neg (wi::to_widest (unit
))
123 == wi::to_widest (step
));
124 *cond_code
= GT_EXPR
;
131 if (POINTER_TYPE_P (TREE_TYPE (*n2
)))
132 *n2
= fold_build_pointer_plus_hwi_loc (loc
, *n2
, 1);
134 *n2
= fold_build2_loc (loc
, PLUS_EXPR
, TREE_TYPE (*n2
), *n2
,
135 build_int_cst (TREE_TYPE (*n2
), 1));
136 *cond_code
= LT_EXPR
;
139 if (POINTER_TYPE_P (TREE_TYPE (*n2
)))
140 *n2
= fold_build_pointer_plus_hwi_loc (loc
, *n2
, -1);
142 *n2
= fold_build2_loc (loc
, MINUS_EXPR
, TREE_TYPE (*n2
), *n2
,
143 build_int_cst (TREE_TYPE (*n2
), 1));
144 *cond_code
= GT_EXPR
;
151 /* Return the looping step from INCR, extracted from the step of a gimple omp
155 omp_get_for_step_from_incr (location_t loc
, tree incr
)
158 switch (TREE_CODE (incr
))
161 step
= TREE_OPERAND (incr
, 1);
163 case POINTER_PLUS_EXPR
:
164 step
= fold_convert (ssizetype
, TREE_OPERAND (incr
, 1));
167 step
= TREE_OPERAND (incr
, 1);
168 step
= fold_build1_loc (loc
, NEGATE_EXPR
, TREE_TYPE (step
), step
);
176 /* Extract the header elements of parallel loop FOR_STMT and store
180 omp_extract_for_data (gomp_for
*for_stmt
, struct omp_for_data
*fd
,
181 struct omp_for_data_loop
*loops
)
183 tree t
, var
, *collapse_iter
, *collapse_count
;
184 tree count
= NULL_TREE
, iter_type
= long_integer_type_node
;
185 struct omp_for_data_loop
*loop
;
187 struct omp_for_data_loop dummy_loop
;
188 location_t loc
= gimple_location (for_stmt
);
189 bool simd
= gimple_omp_for_kind (for_stmt
) == GF_OMP_FOR_KIND_SIMD
;
190 bool distribute
= gimple_omp_for_kind (for_stmt
)
191 == GF_OMP_FOR_KIND_DISTRIBUTE
;
192 bool taskloop
= gimple_omp_for_kind (for_stmt
)
193 == GF_OMP_FOR_KIND_TASKLOOP
;
196 fd
->for_stmt
= for_stmt
;
198 fd
->have_nowait
= distribute
|| simd
;
199 fd
->have_ordered
= false;
200 fd
->have_reductemp
= false;
201 fd
->have_pointer_condtemp
= false;
202 fd
->have_scantemp
= false;
203 fd
->have_nonctrl_scantemp
= false;
204 fd
->lastprivate_conditional
= 0;
205 fd
->tiling
= NULL_TREE
;
208 fd
->sched_kind
= OMP_CLAUSE_SCHEDULE_STATIC
;
209 fd
->sched_modifiers
= 0;
210 fd
->chunk_size
= NULL_TREE
;
211 fd
->simd_schedule
= false;
212 collapse_iter
= NULL
;
213 collapse_count
= NULL
;
215 for (t
= gimple_omp_for_clauses (for_stmt
); t
; t
= OMP_CLAUSE_CHAIN (t
))
216 switch (OMP_CLAUSE_CODE (t
))
218 case OMP_CLAUSE_NOWAIT
:
219 fd
->have_nowait
= true;
221 case OMP_CLAUSE_ORDERED
:
222 fd
->have_ordered
= true;
223 if (OMP_CLAUSE_ORDERED_EXPR (t
))
224 fd
->ordered
= tree_to_shwi (OMP_CLAUSE_ORDERED_EXPR (t
));
226 case OMP_CLAUSE_SCHEDULE
:
227 gcc_assert (!distribute
&& !taskloop
);
229 = (enum omp_clause_schedule_kind
)
230 (OMP_CLAUSE_SCHEDULE_KIND (t
) & OMP_CLAUSE_SCHEDULE_MASK
);
231 fd
->sched_modifiers
= (OMP_CLAUSE_SCHEDULE_KIND (t
)
232 & ~OMP_CLAUSE_SCHEDULE_MASK
);
233 fd
->chunk_size
= OMP_CLAUSE_SCHEDULE_CHUNK_EXPR (t
);
234 fd
->simd_schedule
= OMP_CLAUSE_SCHEDULE_SIMD (t
);
236 case OMP_CLAUSE_DIST_SCHEDULE
:
237 gcc_assert (distribute
);
238 fd
->chunk_size
= OMP_CLAUSE_DIST_SCHEDULE_CHUNK_EXPR (t
);
240 case OMP_CLAUSE_COLLAPSE
:
241 fd
->collapse
= tree_to_shwi (OMP_CLAUSE_COLLAPSE_EXPR (t
));
242 if (fd
->collapse
> 1)
244 collapse_iter
= &OMP_CLAUSE_COLLAPSE_ITERVAR (t
);
245 collapse_count
= &OMP_CLAUSE_COLLAPSE_COUNT (t
);
248 case OMP_CLAUSE_TILE
:
249 fd
->tiling
= OMP_CLAUSE_TILE_LIST (t
);
250 fd
->collapse
= list_length (fd
->tiling
);
251 gcc_assert (fd
->collapse
);
252 collapse_iter
= &OMP_CLAUSE_TILE_ITERVAR (t
);
253 collapse_count
= &OMP_CLAUSE_TILE_COUNT (t
);
255 case OMP_CLAUSE__REDUCTEMP_
:
256 fd
->have_reductemp
= true;
258 case OMP_CLAUSE_LASTPRIVATE
:
259 if (OMP_CLAUSE_LASTPRIVATE_CONDITIONAL (t
))
260 fd
->lastprivate_conditional
++;
262 case OMP_CLAUSE__CONDTEMP_
:
263 if (POINTER_TYPE_P (TREE_TYPE (OMP_CLAUSE_DECL (t
))))
264 fd
->have_pointer_condtemp
= true;
266 case OMP_CLAUSE__SCANTEMP_
:
267 fd
->have_scantemp
= true;
268 if (!OMP_CLAUSE__SCANTEMP__ALLOC (t
)
269 && !OMP_CLAUSE__SCANTEMP__CONTROL (t
))
270 fd
->have_nonctrl_scantemp
= true;
276 if (fd
->collapse
> 1 || fd
->tiling
)
279 fd
->loops
= &fd
->loop
;
281 if (fd
->ordered
&& fd
->collapse
== 1 && loops
!= NULL
)
286 collapse_iter
= &iterv
;
287 collapse_count
= &countv
;
290 /* FIXME: for now map schedule(auto) to schedule(static).
291 There should be analysis to determine whether all iterations
292 are approximately the same amount of work (then schedule(static)
293 is best) or if it varies (then schedule(dynamic,N) is better). */
294 if (fd
->sched_kind
== OMP_CLAUSE_SCHEDULE_AUTO
)
296 fd
->sched_kind
= OMP_CLAUSE_SCHEDULE_STATIC
;
297 gcc_assert (fd
->chunk_size
== NULL
);
299 gcc_assert ((fd
->collapse
== 1 && !fd
->tiling
) || collapse_iter
!= NULL
);
301 fd
->sched_kind
= OMP_CLAUSE_SCHEDULE_RUNTIME
;
302 if (fd
->sched_kind
== OMP_CLAUSE_SCHEDULE_RUNTIME
)
303 gcc_assert (fd
->chunk_size
== NULL
);
304 else if (fd
->chunk_size
== NULL
)
306 /* We only need to compute a default chunk size for ordered
307 static loops and dynamic loops. */
308 if (fd
->sched_kind
!= OMP_CLAUSE_SCHEDULE_STATIC
310 fd
->chunk_size
= (fd
->sched_kind
== OMP_CLAUSE_SCHEDULE_STATIC
)
311 ? integer_zero_node
: integer_one_node
;
314 int cnt
= fd
->ordered
? fd
->ordered
: fd
->collapse
;
315 for (i
= 0; i
< cnt
; i
++)
320 && (fd
->ordered
== 0 || loops
== NULL
))
322 else if (loops
!= NULL
)
327 loop
->v
= gimple_omp_for_index (for_stmt
, i
);
328 gcc_assert (SSA_VAR_P (loop
->v
));
329 gcc_assert (TREE_CODE (TREE_TYPE (loop
->v
)) == INTEGER_TYPE
330 || TREE_CODE (TREE_TYPE (loop
->v
)) == POINTER_TYPE
);
331 var
= TREE_CODE (loop
->v
) == SSA_NAME
? SSA_NAME_VAR (loop
->v
) : loop
->v
;
332 loop
->n1
= gimple_omp_for_initial (for_stmt
, i
);
334 loop
->cond_code
= gimple_omp_for_cond (for_stmt
, i
);
335 loop
->n2
= gimple_omp_for_final (for_stmt
, i
);
336 gcc_assert (loop
->cond_code
!= NE_EXPR
337 || (gimple_omp_for_kind (for_stmt
)
338 != GF_OMP_FOR_KIND_OACC_LOOP
));
340 t
= gimple_omp_for_incr (for_stmt
, i
);
341 gcc_assert (TREE_OPERAND (t
, 0) == var
);
342 loop
->step
= omp_get_for_step_from_incr (loc
, t
);
344 omp_adjust_for_condition (loc
, &loop
->cond_code
, &loop
->n2
, loop
->v
,
348 || (fd
->sched_kind
== OMP_CLAUSE_SCHEDULE_STATIC
349 && !fd
->have_ordered
))
351 if (fd
->collapse
== 1 && !fd
->tiling
)
352 iter_type
= TREE_TYPE (loop
->v
);
354 || TYPE_PRECISION (iter_type
)
355 < TYPE_PRECISION (TREE_TYPE (loop
->v
)))
357 = build_nonstandard_integer_type
358 (TYPE_PRECISION (TREE_TYPE (loop
->v
)), 1);
360 else if (iter_type
!= long_long_unsigned_type_node
)
362 if (POINTER_TYPE_P (TREE_TYPE (loop
->v
)))
363 iter_type
= long_long_unsigned_type_node
;
364 else if (TYPE_UNSIGNED (TREE_TYPE (loop
->v
))
365 && TYPE_PRECISION (TREE_TYPE (loop
->v
))
366 >= TYPE_PRECISION (iter_type
))
370 if (loop
->cond_code
== LT_EXPR
)
371 n
= fold_build2_loc (loc
, PLUS_EXPR
, TREE_TYPE (loop
->v
),
372 loop
->n2
, loop
->step
);
375 if (TREE_CODE (n
) != INTEGER_CST
376 || tree_int_cst_lt (TYPE_MAX_VALUE (iter_type
), n
))
377 iter_type
= long_long_unsigned_type_node
;
379 else if (TYPE_PRECISION (TREE_TYPE (loop
->v
))
380 > TYPE_PRECISION (iter_type
))
384 if (loop
->cond_code
== LT_EXPR
)
387 n2
= fold_build2_loc (loc
, PLUS_EXPR
, TREE_TYPE (loop
->v
),
388 loop
->n2
, loop
->step
);
392 n1
= fold_build2_loc (loc
, MINUS_EXPR
, TREE_TYPE (loop
->v
),
393 loop
->n2
, loop
->step
);
396 if (TREE_CODE (n1
) != INTEGER_CST
397 || TREE_CODE (n2
) != INTEGER_CST
398 || !tree_int_cst_lt (TYPE_MIN_VALUE (iter_type
), n1
)
399 || !tree_int_cst_lt (n2
, TYPE_MAX_VALUE (iter_type
)))
400 iter_type
= long_long_unsigned_type_node
;
404 if (i
>= fd
->collapse
)
407 if (collapse_count
&& *collapse_count
== NULL
)
409 t
= fold_binary (loop
->cond_code
, boolean_type_node
,
410 fold_convert (TREE_TYPE (loop
->v
), loop
->n1
),
411 fold_convert (TREE_TYPE (loop
->v
), loop
->n2
));
412 if (t
&& integer_zerop (t
))
413 count
= build_zero_cst (long_long_unsigned_type_node
);
414 else if ((i
== 0 || count
!= NULL_TREE
)
415 && TREE_CODE (TREE_TYPE (loop
->v
)) == INTEGER_TYPE
416 && TREE_CONSTANT (loop
->n1
)
417 && TREE_CONSTANT (loop
->n2
)
418 && TREE_CODE (loop
->step
) == INTEGER_CST
)
420 tree itype
= TREE_TYPE (loop
->v
);
422 if (POINTER_TYPE_P (itype
))
423 itype
= signed_type_for (itype
);
424 t
= build_int_cst (itype
, (loop
->cond_code
== LT_EXPR
? -1 : 1));
425 t
= fold_build2_loc (loc
, PLUS_EXPR
, itype
,
426 fold_convert_loc (loc
, itype
, loop
->step
),
428 t
= fold_build2_loc (loc
, PLUS_EXPR
, itype
, t
,
429 fold_convert_loc (loc
, itype
, loop
->n2
));
430 t
= fold_build2_loc (loc
, MINUS_EXPR
, itype
, t
,
431 fold_convert_loc (loc
, itype
, loop
->n1
));
432 if (TYPE_UNSIGNED (itype
) && loop
->cond_code
== GT_EXPR
)
434 tree step
= fold_convert_loc (loc
, itype
, loop
->step
);
435 t
= fold_build2_loc (loc
, TRUNC_DIV_EXPR
, itype
,
436 fold_build1_loc (loc
, NEGATE_EXPR
,
438 fold_build1_loc (loc
, NEGATE_EXPR
,
442 t
= fold_build2_loc (loc
, TRUNC_DIV_EXPR
, itype
, t
,
443 fold_convert_loc (loc
, itype
,
445 t
= fold_convert_loc (loc
, long_long_unsigned_type_node
, t
);
446 if (count
!= NULL_TREE
)
447 count
= fold_build2_loc (loc
, MULT_EXPR
,
448 long_long_unsigned_type_node
,
452 if (TREE_CODE (count
) != INTEGER_CST
)
455 else if (count
&& !integer_zerop (count
))
462 && (fd
->sched_kind
!= OMP_CLAUSE_SCHEDULE_STATIC
463 || fd
->have_ordered
))
465 if (!tree_int_cst_lt (count
, TYPE_MAX_VALUE (long_integer_type_node
)))
466 iter_type
= long_long_unsigned_type_node
;
468 iter_type
= long_integer_type_node
;
470 else if (collapse_iter
&& *collapse_iter
!= NULL
)
471 iter_type
= TREE_TYPE (*collapse_iter
);
472 fd
->iter_type
= iter_type
;
473 if (collapse_iter
&& *collapse_iter
== NULL
)
474 *collapse_iter
= create_tmp_var (iter_type
, ".iter");
475 if (collapse_count
&& *collapse_count
== NULL
)
478 *collapse_count
= fold_convert_loc (loc
, iter_type
, count
);
480 *collapse_count
= create_tmp_var (iter_type
, ".count");
483 if (fd
->collapse
> 1 || fd
->tiling
|| (fd
->ordered
&& loops
))
485 fd
->loop
.v
= *collapse_iter
;
486 fd
->loop
.n1
= build_int_cst (TREE_TYPE (fd
->loop
.v
), 0);
487 fd
->loop
.n2
= *collapse_count
;
488 fd
->loop
.step
= build_int_cst (TREE_TYPE (fd
->loop
.v
), 1);
489 fd
->loop
.cond_code
= LT_EXPR
;
495 /* Build a call to GOMP_barrier. */
498 omp_build_barrier (tree lhs
)
500 tree fndecl
= builtin_decl_explicit (lhs
? BUILT_IN_GOMP_BARRIER_CANCEL
501 : BUILT_IN_GOMP_BARRIER
);
502 gcall
*g
= gimple_build_call (fndecl
, 0);
504 gimple_call_set_lhs (g
, lhs
);
508 /* Find OMP_FOR resp. OMP_SIMD with non-NULL OMP_FOR_INIT. Also, fill in pdata
509 array, pdata[0] non-NULL if there is anything non-trivial in between,
510 pdata[1] is address of OMP_PARALLEL in between if any, pdata[2] is address
511 of OMP_FOR in between if any and pdata[3] is address of the inner
515 find_combined_omp_for (tree
*tp
, int *walk_subtrees
, void *data
)
517 tree
**pdata
= (tree
**) data
;
519 switch (TREE_CODE (*tp
))
522 if (OMP_FOR_INIT (*tp
) != NULL_TREE
)
531 if (OMP_FOR_INIT (*tp
) != NULL_TREE
)
538 if (BIND_EXPR_VARS (*tp
)
539 || (BIND_EXPR_BLOCK (*tp
)
540 && BLOCK_VARS (BIND_EXPR_BLOCK (*tp
))))
545 if (!tsi_one_before_end_p (tsi_start (*tp
)))
549 case TRY_FINALLY_EXPR
:
563 /* Return maximum possible vectorization factor for the target. */
570 || !flag_tree_loop_optimize
571 || (!flag_tree_loop_vectorize
572 && global_options_set
.x_flag_tree_loop_vectorize
))
575 auto_vector_modes modes
;
576 targetm
.vectorize
.autovectorize_vector_modes (&modes
, true);
577 if (!modes
.is_empty ())
580 for (unsigned int i
= 0; i
< modes
.length (); ++i
)
581 /* The returned modes use the smallest element size (and thus
582 the largest nunits) for the vectorization approach that they
584 vf
= ordered_max (vf
, GET_MODE_NUNITS (modes
[i
]));
588 machine_mode vqimode
= targetm
.vectorize
.preferred_simd_mode (QImode
);
589 if (GET_MODE_CLASS (vqimode
) == MODE_VECTOR_INT
)
590 return GET_MODE_NUNITS (vqimode
);
595 /* Return maximum SIMT width if offloading may target SIMT hardware. */
598 omp_max_simt_vf (void)
602 if (ENABLE_OFFLOADING
)
603 for (const char *c
= getenv ("OFFLOAD_TARGET_NAMES"); c
;)
605 if (!strncmp (c
, "nvptx", strlen ("nvptx")))
607 else if ((c
= strchr (c
, ':')))
613 /* Store the construct selectors as tree codes from last to first,
614 return their number. */
617 omp_constructor_traits_to_codes (tree ctx
, enum tree_code
*constructs
)
619 int nconstructs
= list_length (ctx
);
620 int i
= nconstructs
- 1;
621 for (tree t2
= ctx
; t2
; t2
= TREE_CHAIN (t2
), i
--)
623 const char *sel
= IDENTIFIER_POINTER (TREE_PURPOSE (t2
));
624 if (!strcmp (sel
, "target"))
625 constructs
[i
] = OMP_TARGET
;
626 else if (!strcmp (sel
, "teams"))
627 constructs
[i
] = OMP_TEAMS
;
628 else if (!strcmp (sel
, "parallel"))
629 constructs
[i
] = OMP_PARALLEL
;
630 else if (!strcmp (sel
, "for") || !strcmp (sel
, "do"))
631 constructs
[i
] = OMP_FOR
;
632 else if (!strcmp (sel
, "simd"))
633 constructs
[i
] = OMP_SIMD
;
637 gcc_assert (i
== -1);
641 /* Return true if PROP is possibly present in one of the offloading target's
642 OpenMP contexts. The format of PROPS string is always offloading target's
643 name terminated by '\0', followed by properties for that offloading
644 target separated by '\0' and terminated by another '\0'. The strings
645 are created from omp-device-properties installed files of all configured
646 offloading targets. */
649 omp_offload_device_kind_arch_isa (const char *props
, const char *prop
)
651 const char *names
= getenv ("OFFLOAD_TARGET_NAMES");
652 if (names
== NULL
|| *names
== '\0')
654 while (*props
!= '\0')
656 size_t name_len
= strlen (props
);
657 bool matches
= false;
658 for (const char *c
= names
; c
; )
660 if (strncmp (props
, c
, name_len
) == 0
661 && (c
[name_len
] == '\0'
662 || c
[name_len
] == ':'
663 || c
[name_len
] == '='))
668 else if ((c
= strchr (c
, ':')))
671 props
= props
+ name_len
+ 1;
672 while (*props
!= '\0')
674 if (matches
&& strcmp (props
, prop
) == 0)
676 props
= strchr (props
, '\0') + 1;
683 /* Return true if the current code location is or might be offloaded.
684 Return true in declare target functions, or when nested in a target
685 region or when unsure, return false otherwise. */
688 omp_maybe_offloaded (void)
690 if (!hsa_gen_requested_p ())
692 if (!ENABLE_OFFLOADING
)
694 const char *names
= getenv ("OFFLOAD_TARGET_NAMES");
695 if (names
== NULL
|| *names
== '\0')
698 if (symtab
->state
== PARSING
)
701 if (cfun
&& cfun
->after_inlining
)
703 if (current_function_decl
704 && lookup_attribute ("omp declare target",
705 DECL_ATTRIBUTES (current_function_decl
)))
707 if (cfun
&& (cfun
->curr_properties
& PROP_gimple_any
) == 0)
709 enum tree_code construct
= OMP_TARGET
;
710 if (omp_construct_selector_matches (&construct
, 1, NULL
))
716 /* Return a name from PROP, a property in selectors accepting
720 omp_context_name_list_prop (tree prop
)
722 if (TREE_PURPOSE (prop
))
723 return IDENTIFIER_POINTER (TREE_PURPOSE (prop
));
726 const char *ret
= TREE_STRING_POINTER (TREE_VALUE (prop
));
727 if ((size_t) TREE_STRING_LENGTH (TREE_VALUE (prop
)) == strlen (ret
) + 1)
733 /* Return 1 if context selector matches the current OpenMP context, 0
734 if it does not and -1 if it is unknown and need to be determined later.
735 Some properties can be checked right away during parsing (this routine),
736 others need to wait until the whole TU is parsed, others need to wait until
737 IPA, others until vectorization. */
740 omp_context_selector_matches (tree ctx
)
743 for (tree t1
= ctx
; t1
; t1
= TREE_CHAIN (t1
))
745 char set
= IDENTIFIER_POINTER (TREE_PURPOSE (t1
))[0];
748 /* For now, ignore the construct set. While something can be
749 determined already during parsing, we don't know until end of TU
750 whether additional constructs aren't added through declare variant
751 unless "omp declare variant variant" attribute exists already
752 (so in most of the cases), and we'd need to maintain set of
753 surrounding OpenMP constructs, which is better handled during
755 if (symtab
->state
== PARSING
)
761 enum tree_code constructs
[5];
763 = omp_constructor_traits_to_codes (TREE_VALUE (t1
), constructs
);
765 if (cfun
&& (cfun
->curr_properties
& PROP_gimple_any
) != 0)
767 if (!cfun
->after_inlining
)
773 for (i
= 0; i
< nconstructs
; ++i
)
774 if (constructs
[i
] == OMP_SIMD
)
781 /* If there is no simd, assume it is ok after IPA,
782 constructs should have been checked before. */
786 int r
= omp_construct_selector_matches (constructs
, nconstructs
,
794 for (tree t2
= TREE_VALUE (t1
); t2
; t2
= TREE_CHAIN (t2
))
796 const char *sel
= IDENTIFIER_POINTER (TREE_PURPOSE (t2
));
800 if (set
== 'i' && !strcmp (sel
, "vendor"))
801 for (tree t3
= TREE_VALUE (t2
); t3
; t3
= TREE_CHAIN (t3
))
803 const char *prop
= omp_context_name_list_prop (t3
);
806 if ((!strcmp (prop
, " score") && TREE_PURPOSE (t3
))
807 || !strcmp (prop
, "gnu"))
813 if (set
== 'i' && !strcmp (sel
, "extension"))
814 /* We don't support any extensions right now. */
818 if (set
== 'i' && !strcmp (sel
, "atomic_default_mem_order"))
820 if (cfun
&& (cfun
->curr_properties
& PROP_gimple_any
) != 0)
823 enum omp_memory_order omo
824 = ((enum omp_memory_order
)
826 & OMP_REQUIRES_ATOMIC_DEFAULT_MEM_ORDER
));
827 if (omo
== OMP_MEMORY_ORDER_UNSPECIFIED
)
829 /* We don't know yet, until end of TU. */
830 if (symtab
->state
== PARSING
)
836 omo
= OMP_MEMORY_ORDER_RELAXED
;
838 tree t3
= TREE_VALUE (t2
);
839 const char *prop
= IDENTIFIER_POINTER (TREE_PURPOSE (t3
));
840 if (!strcmp (prop
, " score"))
842 t3
= TREE_CHAIN (t3
);
843 prop
= IDENTIFIER_POINTER (TREE_PURPOSE (t3
));
845 if (!strcmp (prop
, "relaxed")
846 && omo
!= OMP_MEMORY_ORDER_RELAXED
)
848 else if (!strcmp (prop
, "seq_cst")
849 && omo
!= OMP_MEMORY_ORDER_SEQ_CST
)
851 else if (!strcmp (prop
, "acq_rel")
852 && omo
!= OMP_MEMORY_ORDER_ACQ_REL
)
855 if (set
== 'd' && !strcmp (sel
, "arch"))
856 for (tree t3
= TREE_VALUE (t2
); t3
; t3
= TREE_CHAIN (t3
))
858 const char *arch
= omp_context_name_list_prop (t3
);
862 if (targetm
.omp
.device_kind_arch_isa
!= NULL
)
863 r
= targetm
.omp
.device_kind_arch_isa (omp_device_arch
,
865 if (r
== 0 || (r
== -1 && symtab
->state
!= PARSING
))
867 /* If we are or might be in a target region or
868 declare target function, need to take into account
869 also offloading values. */
870 if (!omp_maybe_offloaded ())
872 if (strcmp (arch
, "hsa") == 0
873 && hsa_gen_requested_p ())
878 if (ENABLE_OFFLOADING
)
880 const char *arches
= omp_offload_device_arch
;
881 if (omp_offload_device_kind_arch_isa (arches
,
892 /* If arch matches on the host, it still might not match
893 in the offloading region. */
894 else if (omp_maybe_offloaded ())
899 if (set
== 'i' && !strcmp (sel
, "unified_address"))
901 if (cfun
&& (cfun
->curr_properties
& PROP_gimple_any
) != 0)
904 if ((omp_requires_mask
& OMP_REQUIRES_UNIFIED_ADDRESS
) == 0)
906 if (symtab
->state
== PARSING
)
913 if (set
== 'i' && !strcmp (sel
, "unified_shared_memory"))
915 if (cfun
&& (cfun
->curr_properties
& PROP_gimple_any
) != 0)
918 if ((omp_requires_mask
919 & OMP_REQUIRES_UNIFIED_SHARED_MEMORY
) == 0)
921 if (symtab
->state
== PARSING
)
930 if (set
== 'i' && !strcmp (sel
, "dynamic_allocators"))
932 if (cfun
&& (cfun
->curr_properties
& PROP_gimple_any
) != 0)
935 if ((omp_requires_mask
936 & OMP_REQUIRES_DYNAMIC_ALLOCATORS
) == 0)
938 if (symtab
->state
== PARSING
)
947 if (set
== 'i' && !strcmp (sel
, "reverse_offload"))
949 if (cfun
&& (cfun
->curr_properties
& PROP_gimple_any
) != 0)
952 if ((omp_requires_mask
& OMP_REQUIRES_REVERSE_OFFLOAD
) == 0)
954 if (symtab
->state
== PARSING
)
963 if (set
== 'd' && !strcmp (sel
, "kind"))
964 for (tree t3
= TREE_VALUE (t2
); t3
; t3
= TREE_CHAIN (t3
))
966 const char *prop
= omp_context_name_list_prop (t3
);
969 if (!strcmp (prop
, "any"))
971 if (!strcmp (prop
, "host"))
973 if (omp_maybe_offloaded ())
977 if (!strcmp (prop
, "nohost"))
979 if (omp_maybe_offloaded ())
986 if (targetm
.omp
.device_kind_arch_isa
!= NULL
)
987 r
= targetm
.omp
.device_kind_arch_isa (omp_device_kind
,
990 r
= strcmp (prop
, "cpu") == 0;
991 if (r
== 0 || (r
== -1 && symtab
->state
!= PARSING
))
993 /* If we are or might be in a target region or
994 declare target function, need to take into account
995 also offloading values. */
996 if (!omp_maybe_offloaded ())
998 if (strcmp (prop
, "gpu") == 0
999 && hsa_gen_requested_p ())
1004 if (ENABLE_OFFLOADING
)
1006 const char *kinds
= omp_offload_device_kind
;
1007 if (omp_offload_device_kind_arch_isa (kinds
, prop
))
1017 /* If kind matches on the host, it still might not match
1018 in the offloading region. */
1019 else if (omp_maybe_offloaded ())
1024 if (set
== 'd' && !strcmp (sel
, "isa"))
1025 for (tree t3
= TREE_VALUE (t2
); t3
; t3
= TREE_CHAIN (t3
))
1027 const char *isa
= omp_context_name_list_prop (t3
);
1031 if (targetm
.omp
.device_kind_arch_isa
!= NULL
)
1032 r
= targetm
.omp
.device_kind_arch_isa (omp_device_isa
,
1034 if (r
== 0 || (r
== -1 && symtab
->state
!= PARSING
))
1036 /* If isa is valid on the target, but not in the
1037 current function and current function has
1038 #pragma omp declare simd on it, some simd clones
1039 might have the isa added later on. */
1041 && targetm
.simd_clone
.compute_vecsize_and_simdlen
1042 && (cfun
== NULL
|| !cfun
->after_inlining
))
1045 = DECL_ATTRIBUTES (current_function_decl
);
1046 if (lookup_attribute ("omp declare simd", attrs
))
1052 /* If we are or might be in a target region or
1053 declare target function, need to take into account
1054 also offloading values. */
1055 if (!omp_maybe_offloaded ())
1057 if (ENABLE_OFFLOADING
)
1059 const char *isas
= omp_offload_device_isa
;
1060 if (omp_offload_device_kind_arch_isa (isas
, isa
))
1070 /* If isa matches on the host, it still might not match
1071 in the offloading region. */
1072 else if (omp_maybe_offloaded ())
1077 if (set
== 'u' && !strcmp (sel
, "condition"))
1078 for (tree t3
= TREE_VALUE (t2
); t3
; t3
= TREE_CHAIN (t3
))
1079 if (TREE_PURPOSE (t3
) == NULL_TREE
)
1081 if (integer_zerop (TREE_VALUE (t3
)))
1083 if (integer_nonzerop (TREE_VALUE (t3
)))
1096 /* Compare construct={simd} CLAUSES1 with CLAUSES2, return 0/-1/1/2 as
1097 in omp_context_selector_set_compare. */
1100 omp_construct_simd_compare (tree clauses1
, tree clauses2
)
1102 if (clauses1
== NULL_TREE
)
1103 return clauses2
== NULL_TREE
? 0 : -1;
1104 if (clauses2
== NULL_TREE
)
1108 struct declare_variant_simd_data
{
1109 bool inbranch
, notinbranch
;
1111 auto_vec
<tree
,16> data_sharing
;
1112 auto_vec
<tree
,16> aligned
;
1113 declare_variant_simd_data ()
1114 : inbranch(false), notinbranch(false), simdlen(NULL_TREE
) {}
1117 for (i
= 0; i
< 2; i
++)
1118 for (tree c
= i
? clauses2
: clauses1
; c
; c
= OMP_CLAUSE_CHAIN (c
))
1121 switch (OMP_CLAUSE_CODE (c
))
1123 case OMP_CLAUSE_INBRANCH
:
1124 data
[i
].inbranch
= true;
1126 case OMP_CLAUSE_NOTINBRANCH
:
1127 data
[i
].notinbranch
= true;
1129 case OMP_CLAUSE_SIMDLEN
:
1130 data
[i
].simdlen
= OMP_CLAUSE_SIMDLEN_EXPR (c
);
1132 case OMP_CLAUSE_UNIFORM
:
1133 case OMP_CLAUSE_LINEAR
:
1134 v
= &data
[i
].data_sharing
;
1136 case OMP_CLAUSE_ALIGNED
:
1137 v
= &data
[i
].aligned
;
1142 unsigned HOST_WIDE_INT argno
= tree_to_uhwi (OMP_CLAUSE_DECL (c
));
1143 if (argno
>= v
->length ())
1144 v
->safe_grow_cleared (argno
+ 1);
1147 /* Here, r is used as a bitmask, 2 is set if CLAUSES1 has something
1148 CLAUSES2 doesn't, 1 is set if CLAUSES2 has something CLAUSES1
1149 doesn't. Thus, r == 3 implies return value 2, r == 1 implies
1150 -1, r == 2 implies 1 and r == 0 implies 0. */
1151 if (data
[0].inbranch
!= data
[1].inbranch
)
1152 r
|= data
[0].inbranch
? 2 : 1;
1153 if (data
[0].notinbranch
!= data
[1].notinbranch
)
1154 r
|= data
[0].notinbranch
? 2 : 1;
1155 if (!simple_cst_equal (data
[0].simdlen
, data
[1].simdlen
))
1157 if (data
[0].simdlen
&& data
[1].simdlen
)
1159 r
|= data
[0].simdlen
? 2 : 1;
1161 if (data
[0].data_sharing
.length () < data
[1].data_sharing
.length ()
1162 || data
[0].aligned
.length () < data
[1].aligned
.length ())
1165 FOR_EACH_VEC_ELT (data
[0].data_sharing
, i
, c1
)
1167 c2
= (i
< data
[1].data_sharing
.length ()
1168 ? data
[1].data_sharing
[i
] : NULL_TREE
);
1169 if ((c1
== NULL_TREE
) != (c2
== NULL_TREE
))
1171 r
|= c1
!= NULL_TREE
? 2 : 1;
1174 if (c1
== NULL_TREE
)
1176 if (OMP_CLAUSE_CODE (c1
) != OMP_CLAUSE_CODE (c2
))
1178 if (OMP_CLAUSE_CODE (c1
) != OMP_CLAUSE_LINEAR
)
1180 if (OMP_CLAUSE_LINEAR_VARIABLE_STRIDE (c1
)
1181 != OMP_CLAUSE_LINEAR_VARIABLE_STRIDE (c2
))
1183 if (OMP_CLAUSE_LINEAR_KIND (c1
) != OMP_CLAUSE_LINEAR_KIND (c2
))
1185 if (!simple_cst_equal (OMP_CLAUSE_LINEAR_STEP (c1
),
1186 OMP_CLAUSE_LINEAR_STEP (c2
)))
1189 FOR_EACH_VEC_ELT (data
[0].aligned
, i
, c1
)
1191 c2
= i
< data
[1].aligned
.length () ? data
[1].aligned
[i
] : NULL_TREE
;
1192 if ((c1
== NULL_TREE
) != (c2
== NULL_TREE
))
1194 r
|= c1
!= NULL_TREE
? 2 : 1;
1197 if (c1
== NULL_TREE
)
1199 if (!simple_cst_equal (OMP_CLAUSE_ALIGNED_ALIGNMENT (c1
),
1200 OMP_CLAUSE_ALIGNED_ALIGNMENT (c2
)))
1209 default: gcc_unreachable ();
1213 /* Compare properties of selectors SEL from SET other than construct.
1214 Return 0/-1/1/2 as in omp_context_selector_set_compare.
1215 Unlike set names or selector names, properties can have duplicates. */
1218 omp_context_selector_props_compare (const char *set
, const char *sel
,
1219 tree ctx1
, tree ctx2
)
1222 for (int pass
= 0; pass
< 2; pass
++)
1223 for (tree t1
= pass
? ctx2
: ctx1
; t1
; t1
= TREE_CHAIN (t1
))
1226 for (t2
= pass
? ctx1
: ctx2
; t2
; t2
= TREE_CHAIN (t2
))
1227 if (TREE_PURPOSE (t1
) == TREE_PURPOSE (t2
))
1229 if (TREE_PURPOSE (t1
) == NULL_TREE
)
1231 if (set
[0] == 'u' && strcmp (sel
, "condition") == 0)
1233 if (integer_zerop (TREE_VALUE (t1
))
1234 != integer_zerop (TREE_VALUE (t2
)))
1238 if (simple_cst_equal (TREE_VALUE (t1
), TREE_VALUE (t2
)))
1241 else if (strcmp (IDENTIFIER_POINTER (TREE_PURPOSE (t1
)),
1244 if (!simple_cst_equal (TREE_VALUE (t1
), TREE_VALUE (t2
)))
1251 else if (TREE_PURPOSE (t1
)
1252 && TREE_PURPOSE (t2
) == NULL_TREE
1253 && TREE_CODE (TREE_VALUE (t2
)) == STRING_CST
)
1255 const char *p1
= omp_context_name_list_prop (t1
);
1256 const char *p2
= omp_context_name_list_prop (t2
);
1258 && strcmp (p1
, p2
) == 0
1259 && strcmp (p1
, " score"))
1262 else if (TREE_PURPOSE (t1
) == NULL_TREE
1263 && TREE_PURPOSE (t2
)
1264 && TREE_CODE (TREE_VALUE (t1
)) == STRING_CST
)
1266 const char *p1
= omp_context_name_list_prop (t1
);
1267 const char *p2
= omp_context_name_list_prop (t2
);
1269 && strcmp (p1
, p2
) == 0
1270 && strcmp (p1
, " score"))
1273 if (t2
== NULL_TREE
)
1275 int r
= pass
? -1 : 1;
1276 if (ret
&& ret
!= r
)
1290 /* Compare single context selector sets CTX1 and CTX2 with SET name.
1291 Return 0 if CTX1 is equal to CTX2,
1292 -1 if CTX1 is a strict subset of CTX2,
1293 1 if CTX2 is a strict subset of CTX1, or
1294 2 if neither context is a subset of another one. */
1297 omp_context_selector_set_compare (const char *set
, tree ctx1
, tree ctx2
)
1299 bool swapped
= false;
1301 int len1
= list_length (ctx1
);
1302 int len2
= list_length (ctx2
);
1307 std::swap (ctx1
, ctx2
);
1308 std::swap (len1
, len2
);
1314 tree simd
= get_identifier ("simd");
1315 /* Handle construct set specially. In this case the order
1316 of the selector matters too. */
1317 for (t1
= ctx1
; t1
; t1
= TREE_CHAIN (t1
))
1318 if (TREE_PURPOSE (t1
) == TREE_PURPOSE (t2
))
1321 if (TREE_PURPOSE (t1
) == simd
)
1322 r
= omp_construct_simd_compare (TREE_VALUE (t1
),
1324 if (r
== 2 || (ret
&& r
&& (ret
< 0) != (r
< 0)))
1328 t2
= TREE_CHAIN (t2
);
1329 if (t2
== NULL_TREE
)
1331 t1
= TREE_CHAIN (t1
);
1339 if (t2
!= NULL_TREE
)
1341 if (t1
!= NULL_TREE
)
1349 return swapped
? -ret
: ret
;
1351 for (tree t1
= ctx1
; t1
; t1
= TREE_CHAIN (t1
))
1354 for (t2
= ctx2
; t2
; t2
= TREE_CHAIN (t2
))
1355 if (TREE_PURPOSE (t1
) == TREE_PURPOSE (t2
))
1357 const char *sel
= IDENTIFIER_POINTER (TREE_PURPOSE (t1
));
1358 int r
= omp_context_selector_props_compare (set
, sel
,
1361 if (r
== 2 || (ret
&& r
&& (ret
< 0) != (r
< 0)))
1368 if (t2
== NULL_TREE
)
1379 return swapped
? -ret
: ret
;
1382 /* Compare whole context selector specification CTX1 and CTX2.
1383 Return 0 if CTX1 is equal to CTX2,
1384 -1 if CTX1 is a strict subset of CTX2,
1385 1 if CTX2 is a strict subset of CTX1, or
1386 2 if neither context is a subset of another one. */
1389 omp_context_selector_compare (tree ctx1
, tree ctx2
)
1391 bool swapped
= false;
1393 int len1
= list_length (ctx1
);
1394 int len2
= list_length (ctx2
);
1399 std::swap (ctx1
, ctx2
);
1400 std::swap (len1
, len2
);
1402 for (tree t1
= ctx1
; t1
; t1
= TREE_CHAIN (t1
))
1405 for (t2
= ctx2
; t2
; t2
= TREE_CHAIN (t2
))
1406 if (TREE_PURPOSE (t1
) == TREE_PURPOSE (t2
))
1408 const char *set
= IDENTIFIER_POINTER (TREE_PURPOSE (t1
));
1409 int r
= omp_context_selector_set_compare (set
, TREE_VALUE (t1
),
1411 if (r
== 2 || (ret
&& r
&& (ret
< 0) != (r
< 0)))
1418 if (t2
== NULL_TREE
)
1429 return swapped
? -ret
: ret
;
1432 /* From context selector CTX, return trait-selector with name SEL in
1433 trait-selector-set with name SET if any, or NULL_TREE if not found.
1434 If SEL is NULL, return the list of trait-selectors in SET. */
1437 omp_get_context_selector (tree ctx
, const char *set
, const char *sel
)
1439 tree setid
= get_identifier (set
);
1440 tree selid
= sel
? get_identifier (sel
) : NULL_TREE
;
1441 for (tree t1
= ctx
; t1
; t1
= TREE_CHAIN (t1
))
1442 if (TREE_PURPOSE (t1
) == setid
)
1445 return TREE_VALUE (t1
);
1446 for (tree t2
= TREE_VALUE (t1
); t2
; t2
= TREE_CHAIN (t2
))
1447 if (TREE_PURPOSE (t2
) == selid
)
1453 /* Compute *SCORE for context selector CTX. Return true if the score
1454 would be different depending on whether it is a declare simd clone or
1455 not. DECLARE_SIMD should be true for the case when it would be
1456 a declare simd clone. */
1459 omp_context_compute_score (tree ctx
, widest_int
*score
, bool declare_simd
)
1461 tree construct
= omp_get_context_selector (ctx
, "construct", NULL
);
1462 bool has_kind
= omp_get_context_selector (ctx
, "device", "kind");
1463 bool has_arch
= omp_get_context_selector (ctx
, "device", "arch");
1464 bool has_isa
= omp_get_context_selector (ctx
, "device", "isa");
1467 for (tree t1
= ctx
; t1
; t1
= TREE_CHAIN (t1
))
1468 if (TREE_VALUE (t1
) != construct
)
1469 for (tree t2
= TREE_VALUE (t1
); t2
; t2
= TREE_CHAIN (t2
))
1470 if (tree t3
= TREE_VALUE (t2
))
1471 if (TREE_PURPOSE (t3
)
1472 && strcmp (IDENTIFIER_POINTER (TREE_PURPOSE (t3
)), " score") == 0
1473 && TREE_CODE (TREE_VALUE (t3
)) == INTEGER_CST
)
1474 *score
+= wi::to_widest (TREE_VALUE (t3
));
1475 if (construct
|| has_kind
|| has_arch
|| has_isa
)
1478 enum tree_code constructs
[5];
1479 int nconstructs
= 0;
1481 nconstructs
= omp_constructor_traits_to_codes (construct
, constructs
);
1482 if (omp_construct_selector_matches (constructs
, nconstructs
, scores
)
1485 int b
= declare_simd
? nconstructs
+ 1 : 0;
1486 if (scores
[b
+ nconstructs
] + 4U < score
->get_precision ())
1488 for (int n
= 0; n
< nconstructs
; ++n
)
1490 if (scores
[b
+ n
] < 0)
1495 *score
+= wi::shifted_mask
<widest_int
> (scores
[b
+ n
], 1, false);
1498 *score
+= wi::shifted_mask
<widest_int
> (scores
[b
+ nconstructs
],
1501 *score
+= wi::shifted_mask
<widest_int
> (scores
[b
+ nconstructs
] + 1,
1504 *score
+= wi::shifted_mask
<widest_int
> (scores
[b
+ nconstructs
] + 2,
1507 else /* FIXME: Implement this. */
1513 /* Class describing a single variant. */
1514 struct GTY(()) omp_declare_variant_entry
{
1515 /* NODE of the variant. */
1516 cgraph_node
*variant
;
1517 /* Score if not in declare simd clone. */
1519 /* Score if in declare simd clone. */
1520 widest_int score_in_declare_simd_clone
;
1521 /* Context selector for the variant. */
1523 /* True if the context selector is known to match already. */
1527 /* Class describing a function with variants. */
1528 struct GTY((for_user
)) omp_declare_variant_base_entry
{
1529 /* NODE of the base function. */
1531 /* NODE of the artificial function created for the deferred variant
1534 /* Vector of the variants. */
1535 vec
<omp_declare_variant_entry
, va_gc
> *variants
;
1538 struct omp_declare_variant_hasher
1539 : ggc_ptr_hash
<omp_declare_variant_base_entry
> {
1540 static hashval_t
hash (omp_declare_variant_base_entry
*);
1541 static bool equal (omp_declare_variant_base_entry
*,
1542 omp_declare_variant_base_entry
*);
1546 omp_declare_variant_hasher::hash (omp_declare_variant_base_entry
*x
)
1548 inchash::hash hstate
;
1549 hstate
.add_int (DECL_UID (x
->base
->decl
));
1550 hstate
.add_int (x
->variants
->length ());
1551 omp_declare_variant_entry
*variant
;
1553 FOR_EACH_VEC_SAFE_ELT (x
->variants
, i
, variant
)
1555 hstate
.add_int (DECL_UID (variant
->variant
->decl
));
1556 hstate
.add_wide_int (variant
->score
);
1557 hstate
.add_wide_int (variant
->score_in_declare_simd_clone
);
1558 hstate
.add_ptr (variant
->ctx
);
1559 hstate
.add_int (variant
->matches
);
1561 return hstate
.end ();
1565 omp_declare_variant_hasher::equal (omp_declare_variant_base_entry
*x
,
1566 omp_declare_variant_base_entry
*y
)
1568 if (x
->base
!= y
->base
1569 || x
->variants
->length () != y
->variants
->length ())
1571 omp_declare_variant_entry
*variant
;
1573 FOR_EACH_VEC_SAFE_ELT (x
->variants
, i
, variant
)
1574 if (variant
->variant
!= (*y
->variants
)[i
].variant
1575 || variant
->score
!= (*y
->variants
)[i
].score
1576 || (variant
->score_in_declare_simd_clone
1577 != (*y
->variants
)[i
].score_in_declare_simd_clone
)
1578 || variant
->ctx
!= (*y
->variants
)[i
].ctx
1579 || variant
->matches
!= (*y
->variants
)[i
].matches
)
1584 static GTY(()) hash_table
<omp_declare_variant_hasher
> *omp_declare_variants
;
1586 struct omp_declare_variant_alt_hasher
1587 : ggc_ptr_hash
<omp_declare_variant_base_entry
> {
1588 static hashval_t
hash (omp_declare_variant_base_entry
*);
1589 static bool equal (omp_declare_variant_base_entry
*,
1590 omp_declare_variant_base_entry
*);
1594 omp_declare_variant_alt_hasher::hash (omp_declare_variant_base_entry
*x
)
1596 return DECL_UID (x
->node
->decl
);
1600 omp_declare_variant_alt_hasher::equal (omp_declare_variant_base_entry
*x
,
1601 omp_declare_variant_base_entry
*y
)
1603 return x
->node
== y
->node
;
1606 static GTY(()) hash_table
<omp_declare_variant_alt_hasher
>
1607 *omp_declare_variant_alt
;
1609 /* Try to resolve declare variant after gimplification. */
1612 omp_resolve_late_declare_variant (tree alt
)
1614 cgraph_node
*node
= cgraph_node::get (alt
);
1615 cgraph_node
*cur_node
= cgraph_node::get (cfun
->decl
);
1617 || !node
->declare_variant_alt
1618 || !cfun
->after_inlining
)
1621 omp_declare_variant_base_entry entry
;
1624 entry
.variants
= NULL
;
1625 omp_declare_variant_base_entry
*entryp
1626 = omp_declare_variant_alt
->find_with_hash (&entry
, DECL_UID (alt
));
1629 omp_declare_variant_entry
*varentry1
, *varentry2
;
1630 auto_vec
<bool, 16> matches
;
1631 unsigned int nmatches
= 0;
1632 FOR_EACH_VEC_SAFE_ELT (entryp
->variants
, i
, varentry1
)
1634 if (varentry1
->matches
)
1636 /* This has been checked to be ok already. */
1637 matches
.safe_push (true);
1641 switch (omp_context_selector_matches (varentry1
->ctx
))
1644 matches
.safe_push (false);
1649 matches
.safe_push (true);
1656 return entryp
->base
->decl
;
1658 /* A context selector that is a strict subset of another context selector
1659 has a score of zero. */
1660 FOR_EACH_VEC_SAFE_ELT (entryp
->variants
, i
, varentry1
)
1664 vec_safe_iterate (entryp
->variants
, j
, &varentry2
); ++j
)
1667 int r
= omp_context_selector_compare (varentry1
->ctx
,
1671 /* ctx1 is a strict subset of ctx2, ignore ctx1. */
1676 /* ctx2 is a strict subset of ctx1, remove ctx2. */
1681 widest_int max_score
= -1;
1683 FOR_EACH_VEC_SAFE_ELT (entryp
->variants
, i
, varentry1
)
1687 = (cur_node
->simdclone
? varentry1
->score_in_declare_simd_clone
1688 : varentry1
->score
);
1689 if (score
> max_score
)
1692 varentry2
= varentry1
;
1695 return varentry2
->variant
->decl
;
1698 /* Hook to adjust hash tables on cgraph_node removal. */
1701 omp_declare_variant_remove_hook (struct cgraph_node
*node
, void *)
1703 if (!node
->declare_variant_alt
)
1706 /* Drop this hash table completely. */
1707 omp_declare_variants
= NULL
;
1708 /* And remove node from the other hash table. */
1709 if (omp_declare_variant_alt
)
1711 omp_declare_variant_base_entry entry
;
1714 entry
.variants
= NULL
;
1715 omp_declare_variant_alt
->remove_elt_with_hash (&entry
,
1716 DECL_UID (node
->decl
));
1720 /* Try to resolve declare variant, return the variant decl if it should
1721 be used instead of base, or base otherwise. */
1724 omp_resolve_declare_variant (tree base
)
1726 tree variant1
= NULL_TREE
, variant2
= NULL_TREE
;
1727 if (cfun
&& (cfun
->curr_properties
& PROP_gimple_any
) != 0)
1728 return omp_resolve_late_declare_variant (base
);
1730 auto_vec
<tree
, 16> variants
;
1731 auto_vec
<bool, 16> defer
;
1732 bool any_deferred
= false;
1733 for (tree attr
= DECL_ATTRIBUTES (base
); attr
; attr
= TREE_CHAIN (attr
))
1735 attr
= lookup_attribute ("omp declare variant base", attr
);
1736 if (attr
== NULL_TREE
)
1738 if (TREE_CODE (TREE_PURPOSE (TREE_VALUE (attr
))) != FUNCTION_DECL
)
1740 cgraph_node
*node
= cgraph_node::get (base
);
1741 /* If this is already a magic decl created by this function,
1742 don't process it again. */
1743 if (node
&& node
->declare_variant_alt
)
1745 switch (omp_context_selector_matches (TREE_VALUE (TREE_VALUE (attr
))))
1748 /* No match, ignore. */
1751 /* Needs to be deferred. */
1752 any_deferred
= true;
1753 variants
.safe_push (attr
);
1754 defer
.safe_push (true);
1757 variants
.safe_push (attr
);
1758 defer
.safe_push (false);
1762 if (variants
.length () == 0)
1767 widest_int max_score1
= 0;
1768 widest_int max_score2
= 0;
1772 omp_declare_variant_base_entry entry
;
1773 entry
.base
= cgraph_node::get_create (base
);
1775 vec_alloc (entry
.variants
, variants
.length ());
1776 FOR_EACH_VEC_ELT (variants
, i
, attr1
)
1781 tree ctx
= TREE_VALUE (TREE_VALUE (attr1
));
1782 need_two
= omp_context_compute_score (ctx
, &score1
, false);
1784 omp_context_compute_score (ctx
, &score2
, true);
1790 max_score1
= score1
;
1791 max_score2
= score2
;
1800 if (max_score1
== score1
)
1801 variant1
= NULL_TREE
;
1802 else if (score1
> max_score1
)
1804 max_score1
= score1
;
1805 variant1
= defer
[i
] ? NULL_TREE
: attr1
;
1807 if (max_score2
== score2
)
1808 variant2
= NULL_TREE
;
1809 else if (score2
> max_score2
)
1811 max_score2
= score2
;
1812 variant2
= defer
[i
] ? NULL_TREE
: attr1
;
1815 omp_declare_variant_entry varentry
;
1817 = cgraph_node::get_create (TREE_PURPOSE (TREE_VALUE (attr1
)));
1818 varentry
.score
= score1
;
1819 varentry
.score_in_declare_simd_clone
= score2
;
1821 varentry
.matches
= !defer
[i
];
1822 entry
.variants
->quick_push (varentry
);
1825 /* If there is a clear winner variant with the score which is not
1826 deferred, verify it is not a strict subset of any other context
1827 selector and if it is not, it is the best alternative no matter
1828 whether the others do or don't match. */
1829 if (variant1
&& variant1
== variant2
)
1831 tree ctx1
= TREE_VALUE (TREE_VALUE (variant1
));
1832 FOR_EACH_VEC_ELT (variants
, i
, attr2
)
1834 if (attr2
== variant1
)
1836 tree ctx2
= TREE_VALUE (TREE_VALUE (attr2
));
1837 int r
= omp_context_selector_compare (ctx1
, ctx2
);
1840 /* The winner is a strict subset of ctx2, can't
1842 variant1
= NULL_TREE
;
1848 vec_free (entry
.variants
);
1849 return TREE_PURPOSE (TREE_VALUE (variant1
));
1853 static struct cgraph_node_hook_list
*node_removal_hook_holder
;
1854 if (!node_removal_hook_holder
)
1855 node_removal_hook_holder
1856 = symtab
->add_cgraph_removal_hook (omp_declare_variant_remove_hook
,
1859 if (omp_declare_variants
== NULL
)
1860 omp_declare_variants
1861 = hash_table
<omp_declare_variant_hasher
>::create_ggc (64);
1862 omp_declare_variant_base_entry
**slot
1863 = omp_declare_variants
->find_slot (&entry
, INSERT
);
1866 vec_free (entry
.variants
);
1867 return (*slot
)->node
->decl
;
1870 *slot
= ggc_cleared_alloc
<omp_declare_variant_base_entry
> ();
1871 (*slot
)->base
= entry
.base
;
1872 (*slot
)->node
= entry
.base
;
1873 (*slot
)->variants
= entry
.variants
;
1874 tree alt
= build_decl (DECL_SOURCE_LOCATION (base
), FUNCTION_DECL
,
1875 DECL_NAME (base
), TREE_TYPE (base
));
1876 DECL_ARTIFICIAL (alt
) = 1;
1877 DECL_IGNORED_P (alt
) = 1;
1878 TREE_STATIC (alt
) = 1;
1879 tree attributes
= DECL_ATTRIBUTES (base
);
1880 if (lookup_attribute ("noipa", attributes
) == NULL
)
1882 attributes
= tree_cons (get_identifier ("noipa"), NULL
, attributes
);
1883 if (lookup_attribute ("noinline", attributes
) == NULL
)
1884 attributes
= tree_cons (get_identifier ("noinline"), NULL
,
1886 if (lookup_attribute ("noclone", attributes
) == NULL
)
1887 attributes
= tree_cons (get_identifier ("noclone"), NULL
,
1889 if (lookup_attribute ("no_icf", attributes
) == NULL
)
1890 attributes
= tree_cons (get_identifier ("no_icf"), NULL
,
1893 DECL_ATTRIBUTES (alt
) = attributes
;
1894 DECL_INITIAL (alt
) = error_mark_node
;
1895 (*slot
)->node
= cgraph_node::create (alt
);
1896 (*slot
)->node
->declare_variant_alt
= 1;
1897 (*slot
)->node
->create_reference (entry
.base
, IPA_REF_ADDR
);
1898 omp_declare_variant_entry
*varentry
;
1899 FOR_EACH_VEC_SAFE_ELT (entry
.variants
, i
, varentry
)
1900 (*slot
)->node
->create_reference (varentry
->variant
, IPA_REF_ADDR
);
1901 if (omp_declare_variant_alt
== NULL
)
1902 omp_declare_variant_alt
1903 = hash_table
<omp_declare_variant_alt_hasher
>::create_ggc (64);
1904 *omp_declare_variant_alt
->find_slot_with_hash (*slot
, DECL_UID (alt
),
1909 if (variants
.length () == 1)
1910 return TREE_PURPOSE (TREE_VALUE (variants
[0]));
1912 /* A context selector that is a strict subset of another context selector
1913 has a score of zero. */
1916 FOR_EACH_VEC_ELT (variants
, i
, attr1
)
1919 tree ctx1
= TREE_VALUE (TREE_VALUE (attr1
));
1920 FOR_EACH_VEC_ELT_FROM (variants
, j
, attr2
, i
+ 1)
1923 tree ctx2
= TREE_VALUE (TREE_VALUE (attr2
));
1924 int r
= omp_context_selector_compare (ctx1
, ctx2
);
1927 /* ctx1 is a strict subset of ctx2, remove
1928 attr1 from the vector. */
1929 variants
[i
] = NULL_TREE
;
1933 /* ctx2 is a strict subset of ctx1, remove attr2
1935 variants
[j
] = NULL_TREE
;
1938 widest_int max_score1
= 0;
1939 widest_int max_score2
= 0;
1941 FOR_EACH_VEC_ELT (variants
, i
, attr1
)
1953 ctx
= TREE_VALUE (TREE_VALUE (variant1
));
1954 need_two
= omp_context_compute_score (ctx
, &max_score1
, false);
1956 omp_context_compute_score (ctx
, &max_score2
, true);
1958 max_score2
= max_score1
;
1960 ctx
= TREE_VALUE (TREE_VALUE (attr1
));
1961 need_two
= omp_context_compute_score (ctx
, &score1
, false);
1963 omp_context_compute_score (ctx
, &score2
, true);
1966 if (score1
> max_score1
)
1968 max_score1
= score1
;
1971 if (score2
> max_score2
)
1973 max_score2
= score2
;
1983 /* If there is a disagreement on which variant has the highest score
1984 depending on whether it will be in a declare simd clone or not,
1985 punt for now and defer until after IPA where we will know that. */
1986 return ((variant1
&& variant1
== variant2
)
1987 ? TREE_PURPOSE (TREE_VALUE (variant1
)) : base
);
1991 /* Encode an oacc launch argument. This matches the GOMP_LAUNCH_PACK
1992 macro on gomp-constants.h. We do not check for overflow. */
1995 oacc_launch_pack (unsigned code
, tree device
, unsigned op
)
1999 res
= build_int_cst (unsigned_type_node
, GOMP_LAUNCH_PACK (code
, 0, op
));
2002 device
= fold_build2 (LSHIFT_EXPR
, unsigned_type_node
,
2003 device
, build_int_cst (unsigned_type_node
,
2004 GOMP_LAUNCH_DEVICE_SHIFT
));
2005 res
= fold_build2 (BIT_IOR_EXPR
, unsigned_type_node
, res
, device
);
2010 /* FIXME: What is the following comment for? */
2011 /* Look for compute grid dimension clauses and convert to an attribute
2012 attached to FN. This permits the target-side code to (a) massage
2013 the dimensions, (b) emit that data and (c) optimize. Non-constant
2014 dimensions are pushed onto ARGS.
2016 The attribute value is a TREE_LIST. A set of dimensions is
2017 represented as a list of INTEGER_CST. Those that are runtime
2018 exprs are represented as an INTEGER_CST of zero.
2020 TODO: Normally the attribute will just contain a single such list. If
2021 however it contains a list of lists, this will represent the use of
2022 device_type. Each member of the outer list is an assoc list of
2023 dimensions, keyed by the device type. The first entry will be the
2024 default. Well, that's the plan. */
2026 /* Replace any existing oacc fn attribute with updated dimensions. */
2028 /* Variant working on a list of attributes. */
2031 oacc_replace_fn_attrib_attr (tree attribs
, tree dims
)
2033 tree ident
= get_identifier (OACC_FN_ATTRIB
);
2035 /* If we happen to be present as the first attrib, drop it. */
2036 if (attribs
&& TREE_PURPOSE (attribs
) == ident
)
2037 attribs
= TREE_CHAIN (attribs
);
2038 return tree_cons (ident
, dims
, attribs
);
2041 /* Variant working on a function decl. */
2044 oacc_replace_fn_attrib (tree fn
, tree dims
)
2046 DECL_ATTRIBUTES (fn
)
2047 = oacc_replace_fn_attrib_attr (DECL_ATTRIBUTES (fn
), dims
);
2050 /* Scan CLAUSES for launch dimensions and attach them to the oacc
2051 function attribute. Push any that are non-constant onto the ARGS
2052 list, along with an appropriate GOMP_LAUNCH_DIM tag. */
2055 oacc_set_fn_attrib (tree fn
, tree clauses
, vec
<tree
> *args
)
2057 /* Must match GOMP_DIM ordering. */
2058 static const omp_clause_code ids
[]
2059 = { OMP_CLAUSE_NUM_GANGS
, OMP_CLAUSE_NUM_WORKERS
,
2060 OMP_CLAUSE_VECTOR_LENGTH
};
2062 tree dims
[GOMP_DIM_MAX
];
2064 tree attr
= NULL_TREE
;
2065 unsigned non_const
= 0;
2067 for (ix
= GOMP_DIM_MAX
; ix
--;)
2069 tree clause
= omp_find_clause (clauses
, ids
[ix
]);
2070 tree dim
= NULL_TREE
;
2073 dim
= OMP_CLAUSE_EXPR (clause
, ids
[ix
]);
2075 if (dim
&& TREE_CODE (dim
) != INTEGER_CST
)
2077 dim
= integer_zero_node
;
2078 non_const
|= GOMP_DIM_MASK (ix
);
2080 attr
= tree_cons (NULL_TREE
, dim
, attr
);
2083 oacc_replace_fn_attrib (fn
, attr
);
2087 /* Push a dynamic argument set. */
2088 args
->safe_push (oacc_launch_pack (GOMP_LAUNCH_DIM
,
2089 NULL_TREE
, non_const
));
2090 for (unsigned ix
= 0; ix
!= GOMP_DIM_MAX
; ix
++)
2091 if (non_const
& GOMP_DIM_MASK (ix
))
2092 args
->safe_push (dims
[ix
]);
2096 /* Verify OpenACC routine clauses.
2098 Returns 0 if FNDECL should be marked with an OpenACC 'routine' directive, 1
2099 if it has already been marked in compatible way, and -1 if incompatible.
2100 Upon returning, the chain of clauses will contain exactly one clause
2101 specifying the level of parallelism. */
2104 oacc_verify_routine_clauses (tree fndecl
, tree
*clauses
, location_t loc
,
2105 const char *routine_str
)
2107 tree c_level
= NULL_TREE
;
2108 tree c_p
= NULL_TREE
;
2109 for (tree c
= *clauses
; c
; c_p
= c
, c
= OMP_CLAUSE_CHAIN (c
))
2110 switch (OMP_CLAUSE_CODE (c
))
2112 case OMP_CLAUSE_GANG
:
2113 case OMP_CLAUSE_WORKER
:
2114 case OMP_CLAUSE_VECTOR
:
2115 case OMP_CLAUSE_SEQ
:
2116 if (c_level
== NULL_TREE
)
2118 else if (OMP_CLAUSE_CODE (c
) == OMP_CLAUSE_CODE (c_level
))
2120 /* This has already been diagnosed in the front ends. */
2121 /* Drop the duplicate clause. */
2122 gcc_checking_assert (c_p
!= NULL_TREE
);
2123 OMP_CLAUSE_CHAIN (c_p
) = OMP_CLAUSE_CHAIN (c
);
2128 error_at (OMP_CLAUSE_LOCATION (c
),
2129 "%qs specifies a conflicting level of parallelism",
2130 omp_clause_code_name
[OMP_CLAUSE_CODE (c
)]);
2131 inform (OMP_CLAUSE_LOCATION (c_level
),
2132 "... to the previous %qs clause here",
2133 omp_clause_code_name
[OMP_CLAUSE_CODE (c_level
)]);
2134 /* Drop the conflicting clause. */
2135 gcc_checking_assert (c_p
!= NULL_TREE
);
2136 OMP_CLAUSE_CHAIN (c_p
) = OMP_CLAUSE_CHAIN (c
);
2143 if (c_level
== NULL_TREE
)
2145 /* Default to an implicit 'seq' clause. */
2146 c_level
= build_omp_clause (loc
, OMP_CLAUSE_SEQ
);
2147 OMP_CLAUSE_CHAIN (c_level
) = *clauses
;
2150 /* In *clauses, we now have exactly one clause specifying the level of
2154 = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (fndecl
));
2155 if (attr
!= NULL_TREE
)
2157 /* Diagnose if "#pragma omp declare target" has also been applied. */
2158 if (TREE_VALUE (attr
) == NULL_TREE
)
2160 /* See <https://gcc.gnu.org/PR93465>; the semantics of combining
2161 OpenACC and OpenMP 'target' are not clear. */
2163 "cannot apply %<%s%> to %qD, which has also been"
2164 " marked with an OpenMP 'declare target' directive",
2165 routine_str
, fndecl
);
2170 /* If a "#pragma acc routine" has already been applied, just verify
2171 this one for compatibility. */
2172 /* Collect previous directive's clauses. */
2173 tree c_level_p
= NULL_TREE
;
2174 for (tree c
= TREE_VALUE (attr
); c
; c
= OMP_CLAUSE_CHAIN (c
))
2175 switch (OMP_CLAUSE_CODE (c
))
2177 case OMP_CLAUSE_GANG
:
2178 case OMP_CLAUSE_WORKER
:
2179 case OMP_CLAUSE_VECTOR
:
2180 case OMP_CLAUSE_SEQ
:
2181 gcc_checking_assert (c_level_p
== NULL_TREE
);
2187 gcc_checking_assert (c_level_p
!= NULL_TREE
);
2188 /* ..., and compare to current directive's, which we've already collected
2192 /* Matching level of parallelism? */
2193 if (OMP_CLAUSE_CODE (c_level
) != OMP_CLAUSE_CODE (c_level_p
))
2196 c_diag_p
= c_level_p
;
2203 if (c_diag
!= NULL_TREE
)
2204 error_at (OMP_CLAUSE_LOCATION (c_diag
),
2205 "incompatible %qs clause when applying"
2206 " %<%s%> to %qD, which has already been"
2207 " marked with an OpenACC 'routine' directive",
2208 omp_clause_code_name
[OMP_CLAUSE_CODE (c_diag
)],
2209 routine_str
, fndecl
);
2210 else if (c_diag_p
!= NULL_TREE
)
2212 "missing %qs clause when applying"
2213 " %<%s%> to %qD, which has already been"
2214 " marked with an OpenACC 'routine' directive",
2215 omp_clause_code_name
[OMP_CLAUSE_CODE (c_diag_p
)],
2216 routine_str
, fndecl
);
2219 if (c_diag_p
!= NULL_TREE
)
2220 inform (OMP_CLAUSE_LOCATION (c_diag_p
),
2221 "... with %qs clause here",
2222 omp_clause_code_name
[OMP_CLAUSE_CODE (c_diag_p
)]);
2225 /* In the front ends, we don't preserve location information for the
2226 OpenACC routine directive itself. However, that of c_level_p
2228 location_t loc_routine
= OMP_CLAUSE_LOCATION (c_level_p
);
2229 inform (loc_routine
, "... without %qs clause near to here",
2230 omp_clause_code_name
[OMP_CLAUSE_CODE (c_diag
)]);
2239 /* Process the OpenACC 'routine' directive clauses to generate an attribute
2240 for the level of parallelism. All dimensions have a size of zero
2241 (dynamic). TREE_PURPOSE is set to indicate whether that dimension
2242 can have a loop partitioned on it. non-zero indicates
2243 yes, zero indicates no. By construction once a non-zero has been
2244 reached, further inner dimensions must also be non-zero. We set
2245 TREE_VALUE to zero for the dimensions that may be partitioned and
2246 1 for the other ones -- if a loop is (erroneously) spawned at
2247 an outer level, we don't want to try and partition it. */
2250 oacc_build_routine_dims (tree clauses
)
2252 /* Must match GOMP_DIM ordering. */
2253 static const omp_clause_code ids
[]
2254 = {OMP_CLAUSE_GANG
, OMP_CLAUSE_WORKER
, OMP_CLAUSE_VECTOR
, OMP_CLAUSE_SEQ
};
2258 for (; clauses
; clauses
= OMP_CLAUSE_CHAIN (clauses
))
2259 for (ix
= GOMP_DIM_MAX
+ 1; ix
--;)
2260 if (OMP_CLAUSE_CODE (clauses
) == ids
[ix
])
2265 gcc_checking_assert (level
>= 0);
2267 tree dims
= NULL_TREE
;
2269 for (ix
= GOMP_DIM_MAX
; ix
--;)
2270 dims
= tree_cons (build_int_cst (boolean_type_node
, ix
>= level
),
2271 build_int_cst (integer_type_node
, ix
< level
), dims
);
2276 /* Retrieve the oacc function attrib and return it. Non-oacc
2277 functions will return NULL. */
2280 oacc_get_fn_attrib (tree fn
)
2282 return lookup_attribute (OACC_FN_ATTRIB
, DECL_ATTRIBUTES (fn
));
2285 /* Return true if FN is an OpenMP or OpenACC offloading function. */
2288 offloading_function_p (tree fn
)
2290 tree attrs
= DECL_ATTRIBUTES (fn
);
2291 return (lookup_attribute ("omp declare target", attrs
)
2292 || lookup_attribute ("omp target entrypoint", attrs
));
2295 /* Extract an oacc execution dimension from FN. FN must be an
2296 offloaded function or routine that has already had its execution
2297 dimensions lowered to the target-specific values. */
2300 oacc_get_fn_dim_size (tree fn
, int axis
)
2302 tree attrs
= oacc_get_fn_attrib (fn
);
2304 gcc_assert (axis
< GOMP_DIM_MAX
);
2306 tree dims
= TREE_VALUE (attrs
);
2308 dims
= TREE_CHAIN (dims
);
2310 int size
= TREE_INT_CST_LOW (TREE_VALUE (dims
));
2315 /* Extract the dimension axis from an IFN_GOACC_DIM_POS or
2316 IFN_GOACC_DIM_SIZE call. */
2319 oacc_get_ifn_dim_arg (const gimple
*stmt
)
2321 gcc_checking_assert (gimple_call_internal_fn (stmt
) == IFN_GOACC_DIM_SIZE
2322 || gimple_call_internal_fn (stmt
) == IFN_GOACC_DIM_POS
);
2323 tree arg
= gimple_call_arg (stmt
, 0);
2324 HOST_WIDE_INT axis
= TREE_INT_CST_LOW (arg
);
2326 gcc_checking_assert (axis
>= 0 && axis
< GOMP_DIM_MAX
);
2330 #include "gt-omp-general.h"