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-2019 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 "symbol-summary.h"
41 #include "hsa-common.h"
42 #include "tree-pass.h"
44 enum omp_requires omp_requires_mask
;
47 omp_find_clause (tree clauses
, enum omp_clause_code kind
)
49 for (; clauses
; clauses
= OMP_CLAUSE_CHAIN (clauses
))
50 if (OMP_CLAUSE_CODE (clauses
) == kind
)
56 /* True if OpenMP should regard this DECL as being a scalar which has Fortran's
57 allocatable or pointer attribute. */
59 omp_is_allocatable_or_ptr (tree decl
)
61 return lang_hooks
.decls
.omp_is_allocatable_or_ptr (decl
);
64 /* Return true if DECL is a Fortran optional argument. */
67 omp_is_optional_argument (tree decl
)
69 return lang_hooks
.decls
.omp_is_optional_argument (decl
);
72 /* Return true if DECL is a reference type. */
75 omp_is_reference (tree decl
)
77 return lang_hooks
.decls
.omp_privatize_by_reference (decl
);
80 /* Adjust *COND_CODE and *N2 so that the former is either LT_EXPR or GT_EXPR,
81 given that V is the loop index variable and STEP is loop step. */
84 omp_adjust_for_condition (location_t loc
, enum tree_code
*cond_code
, tree
*n2
,
94 gcc_assert (TREE_CODE (step
) == INTEGER_CST
);
95 if (TREE_CODE (TREE_TYPE (v
)) == INTEGER_TYPE
)
97 if (integer_onep (step
))
101 gcc_assert (integer_minus_onep (step
));
102 *cond_code
= GT_EXPR
;
107 tree unit
= TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (v
)));
108 gcc_assert (TREE_CODE (unit
) == INTEGER_CST
);
109 if (tree_int_cst_equal (unit
, step
))
110 *cond_code
= LT_EXPR
;
113 gcc_assert (wi::neg (wi::to_widest (unit
))
114 == wi::to_widest (step
));
115 *cond_code
= GT_EXPR
;
122 if (POINTER_TYPE_P (TREE_TYPE (*n2
)))
123 *n2
= fold_build_pointer_plus_hwi_loc (loc
, *n2
, 1);
125 *n2
= fold_build2_loc (loc
, PLUS_EXPR
, TREE_TYPE (*n2
), *n2
,
126 build_int_cst (TREE_TYPE (*n2
), 1));
127 *cond_code
= LT_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
, MINUS_EXPR
, TREE_TYPE (*n2
), *n2
,
134 build_int_cst (TREE_TYPE (*n2
), 1));
135 *cond_code
= GT_EXPR
;
142 /* Return the looping step from INCR, extracted from the step of a gimple omp
146 omp_get_for_step_from_incr (location_t loc
, tree incr
)
149 switch (TREE_CODE (incr
))
152 step
= TREE_OPERAND (incr
, 1);
154 case POINTER_PLUS_EXPR
:
155 step
= fold_convert (ssizetype
, TREE_OPERAND (incr
, 1));
158 step
= TREE_OPERAND (incr
, 1);
159 step
= fold_build1_loc (loc
, NEGATE_EXPR
, TREE_TYPE (step
), step
);
167 /* Extract the header elements of parallel loop FOR_STMT and store
171 omp_extract_for_data (gomp_for
*for_stmt
, struct omp_for_data
*fd
,
172 struct omp_for_data_loop
*loops
)
174 tree t
, var
, *collapse_iter
, *collapse_count
;
175 tree count
= NULL_TREE
, iter_type
= long_integer_type_node
;
176 struct omp_for_data_loop
*loop
;
178 struct omp_for_data_loop dummy_loop
;
179 location_t loc
= gimple_location (for_stmt
);
180 bool simd
= gimple_omp_for_kind (for_stmt
) == GF_OMP_FOR_KIND_SIMD
;
181 bool distribute
= gimple_omp_for_kind (for_stmt
)
182 == GF_OMP_FOR_KIND_DISTRIBUTE
;
183 bool taskloop
= gimple_omp_for_kind (for_stmt
)
184 == GF_OMP_FOR_KIND_TASKLOOP
;
187 fd
->for_stmt
= for_stmt
;
189 fd
->have_nowait
= distribute
|| simd
;
190 fd
->have_ordered
= false;
191 fd
->have_reductemp
= false;
192 fd
->have_pointer_condtemp
= false;
193 fd
->have_scantemp
= false;
194 fd
->have_nonctrl_scantemp
= false;
195 fd
->lastprivate_conditional
= 0;
196 fd
->tiling
= NULL_TREE
;
199 fd
->sched_kind
= OMP_CLAUSE_SCHEDULE_STATIC
;
200 fd
->sched_modifiers
= 0;
201 fd
->chunk_size
= NULL_TREE
;
202 fd
->simd_schedule
= false;
203 collapse_iter
= NULL
;
204 collapse_count
= NULL
;
206 for (t
= gimple_omp_for_clauses (for_stmt
); t
; t
= OMP_CLAUSE_CHAIN (t
))
207 switch (OMP_CLAUSE_CODE (t
))
209 case OMP_CLAUSE_NOWAIT
:
210 fd
->have_nowait
= true;
212 case OMP_CLAUSE_ORDERED
:
213 fd
->have_ordered
= true;
214 if (OMP_CLAUSE_ORDERED_EXPR (t
))
215 fd
->ordered
= tree_to_shwi (OMP_CLAUSE_ORDERED_EXPR (t
));
217 case OMP_CLAUSE_SCHEDULE
:
218 gcc_assert (!distribute
&& !taskloop
);
220 = (enum omp_clause_schedule_kind
)
221 (OMP_CLAUSE_SCHEDULE_KIND (t
) & OMP_CLAUSE_SCHEDULE_MASK
);
222 fd
->sched_modifiers
= (OMP_CLAUSE_SCHEDULE_KIND (t
)
223 & ~OMP_CLAUSE_SCHEDULE_MASK
);
224 fd
->chunk_size
= OMP_CLAUSE_SCHEDULE_CHUNK_EXPR (t
);
225 fd
->simd_schedule
= OMP_CLAUSE_SCHEDULE_SIMD (t
);
227 case OMP_CLAUSE_DIST_SCHEDULE
:
228 gcc_assert (distribute
);
229 fd
->chunk_size
= OMP_CLAUSE_DIST_SCHEDULE_CHUNK_EXPR (t
);
231 case OMP_CLAUSE_COLLAPSE
:
232 fd
->collapse
= tree_to_shwi (OMP_CLAUSE_COLLAPSE_EXPR (t
));
233 if (fd
->collapse
> 1)
235 collapse_iter
= &OMP_CLAUSE_COLLAPSE_ITERVAR (t
);
236 collapse_count
= &OMP_CLAUSE_COLLAPSE_COUNT (t
);
239 case OMP_CLAUSE_TILE
:
240 fd
->tiling
= OMP_CLAUSE_TILE_LIST (t
);
241 fd
->collapse
= list_length (fd
->tiling
);
242 gcc_assert (fd
->collapse
);
243 collapse_iter
= &OMP_CLAUSE_TILE_ITERVAR (t
);
244 collapse_count
= &OMP_CLAUSE_TILE_COUNT (t
);
246 case OMP_CLAUSE__REDUCTEMP_
:
247 fd
->have_reductemp
= true;
249 case OMP_CLAUSE_LASTPRIVATE
:
250 if (OMP_CLAUSE_LASTPRIVATE_CONDITIONAL (t
))
251 fd
->lastprivate_conditional
++;
253 case OMP_CLAUSE__CONDTEMP_
:
254 if (POINTER_TYPE_P (TREE_TYPE (OMP_CLAUSE_DECL (t
))))
255 fd
->have_pointer_condtemp
= true;
257 case OMP_CLAUSE__SCANTEMP_
:
258 fd
->have_scantemp
= true;
259 if (!OMP_CLAUSE__SCANTEMP__ALLOC (t
)
260 && !OMP_CLAUSE__SCANTEMP__CONTROL (t
))
261 fd
->have_nonctrl_scantemp
= true;
267 if (fd
->collapse
> 1 || fd
->tiling
)
270 fd
->loops
= &fd
->loop
;
272 if (fd
->ordered
&& fd
->collapse
== 1 && loops
!= NULL
)
277 collapse_iter
= &iterv
;
278 collapse_count
= &countv
;
281 /* FIXME: for now map schedule(auto) to schedule(static).
282 There should be analysis to determine whether all iterations
283 are approximately the same amount of work (then schedule(static)
284 is best) or if it varies (then schedule(dynamic,N) is better). */
285 if (fd
->sched_kind
== OMP_CLAUSE_SCHEDULE_AUTO
)
287 fd
->sched_kind
= OMP_CLAUSE_SCHEDULE_STATIC
;
288 gcc_assert (fd
->chunk_size
== NULL
);
290 gcc_assert ((fd
->collapse
== 1 && !fd
->tiling
) || collapse_iter
!= NULL
);
292 fd
->sched_kind
= OMP_CLAUSE_SCHEDULE_RUNTIME
;
293 if (fd
->sched_kind
== OMP_CLAUSE_SCHEDULE_RUNTIME
)
294 gcc_assert (fd
->chunk_size
== NULL
);
295 else if (fd
->chunk_size
== NULL
)
297 /* We only need to compute a default chunk size for ordered
298 static loops and dynamic loops. */
299 if (fd
->sched_kind
!= OMP_CLAUSE_SCHEDULE_STATIC
301 fd
->chunk_size
= (fd
->sched_kind
== OMP_CLAUSE_SCHEDULE_STATIC
)
302 ? integer_zero_node
: integer_one_node
;
305 int cnt
= fd
->ordered
? fd
->ordered
: fd
->collapse
;
306 for (i
= 0; i
< cnt
; i
++)
311 && (fd
->ordered
== 0 || loops
== NULL
))
313 else if (loops
!= NULL
)
318 loop
->v
= gimple_omp_for_index (for_stmt
, i
);
319 gcc_assert (SSA_VAR_P (loop
->v
));
320 gcc_assert (TREE_CODE (TREE_TYPE (loop
->v
)) == INTEGER_TYPE
321 || TREE_CODE (TREE_TYPE (loop
->v
)) == POINTER_TYPE
);
322 var
= TREE_CODE (loop
->v
) == SSA_NAME
? SSA_NAME_VAR (loop
->v
) : loop
->v
;
323 loop
->n1
= gimple_omp_for_initial (for_stmt
, i
);
325 loop
->cond_code
= gimple_omp_for_cond (for_stmt
, i
);
326 loop
->n2
= gimple_omp_for_final (for_stmt
, i
);
327 gcc_assert (loop
->cond_code
!= NE_EXPR
328 || (gimple_omp_for_kind (for_stmt
)
329 != GF_OMP_FOR_KIND_OACC_LOOP
));
331 t
= gimple_omp_for_incr (for_stmt
, i
);
332 gcc_assert (TREE_OPERAND (t
, 0) == var
);
333 loop
->step
= omp_get_for_step_from_incr (loc
, t
);
335 omp_adjust_for_condition (loc
, &loop
->cond_code
, &loop
->n2
, loop
->v
,
339 || (fd
->sched_kind
== OMP_CLAUSE_SCHEDULE_STATIC
340 && !fd
->have_ordered
))
342 if (fd
->collapse
== 1 && !fd
->tiling
)
343 iter_type
= TREE_TYPE (loop
->v
);
345 || TYPE_PRECISION (iter_type
)
346 < TYPE_PRECISION (TREE_TYPE (loop
->v
)))
348 = build_nonstandard_integer_type
349 (TYPE_PRECISION (TREE_TYPE (loop
->v
)), 1);
351 else if (iter_type
!= long_long_unsigned_type_node
)
353 if (POINTER_TYPE_P (TREE_TYPE (loop
->v
)))
354 iter_type
= long_long_unsigned_type_node
;
355 else if (TYPE_UNSIGNED (TREE_TYPE (loop
->v
))
356 && TYPE_PRECISION (TREE_TYPE (loop
->v
))
357 >= TYPE_PRECISION (iter_type
))
361 if (loop
->cond_code
== LT_EXPR
)
362 n
= fold_build2_loc (loc
, PLUS_EXPR
, TREE_TYPE (loop
->v
),
363 loop
->n2
, loop
->step
);
366 if (TREE_CODE (n
) != INTEGER_CST
367 || tree_int_cst_lt (TYPE_MAX_VALUE (iter_type
), n
))
368 iter_type
= long_long_unsigned_type_node
;
370 else if (TYPE_PRECISION (TREE_TYPE (loop
->v
))
371 > TYPE_PRECISION (iter_type
))
375 if (loop
->cond_code
== LT_EXPR
)
378 n2
= fold_build2_loc (loc
, PLUS_EXPR
, TREE_TYPE (loop
->v
),
379 loop
->n2
, loop
->step
);
383 n1
= fold_build2_loc (loc
, MINUS_EXPR
, TREE_TYPE (loop
->v
),
384 loop
->n2
, loop
->step
);
387 if (TREE_CODE (n1
) != INTEGER_CST
388 || TREE_CODE (n2
) != INTEGER_CST
389 || !tree_int_cst_lt (TYPE_MIN_VALUE (iter_type
), n1
)
390 || !tree_int_cst_lt (n2
, TYPE_MAX_VALUE (iter_type
)))
391 iter_type
= long_long_unsigned_type_node
;
395 if (i
>= fd
->collapse
)
398 if (collapse_count
&& *collapse_count
== NULL
)
400 t
= fold_binary (loop
->cond_code
, boolean_type_node
,
401 fold_convert (TREE_TYPE (loop
->v
), loop
->n1
),
402 fold_convert (TREE_TYPE (loop
->v
), loop
->n2
));
403 if (t
&& integer_zerop (t
))
404 count
= build_zero_cst (long_long_unsigned_type_node
);
405 else if ((i
== 0 || count
!= NULL_TREE
)
406 && TREE_CODE (TREE_TYPE (loop
->v
)) == INTEGER_TYPE
407 && TREE_CONSTANT (loop
->n1
)
408 && TREE_CONSTANT (loop
->n2
)
409 && TREE_CODE (loop
->step
) == INTEGER_CST
)
411 tree itype
= TREE_TYPE (loop
->v
);
413 if (POINTER_TYPE_P (itype
))
414 itype
= signed_type_for (itype
);
415 t
= build_int_cst (itype
, (loop
->cond_code
== LT_EXPR
? -1 : 1));
416 t
= fold_build2_loc (loc
, PLUS_EXPR
, itype
,
417 fold_convert_loc (loc
, itype
, loop
->step
),
419 t
= fold_build2_loc (loc
, PLUS_EXPR
, itype
, t
,
420 fold_convert_loc (loc
, itype
, loop
->n2
));
421 t
= fold_build2_loc (loc
, MINUS_EXPR
, itype
, t
,
422 fold_convert_loc (loc
, itype
, loop
->n1
));
423 if (TYPE_UNSIGNED (itype
) && loop
->cond_code
== GT_EXPR
)
425 tree step
= fold_convert_loc (loc
, itype
, loop
->step
);
426 t
= fold_build2_loc (loc
, TRUNC_DIV_EXPR
, itype
,
427 fold_build1_loc (loc
, NEGATE_EXPR
,
429 fold_build1_loc (loc
, NEGATE_EXPR
,
433 t
= fold_build2_loc (loc
, TRUNC_DIV_EXPR
, itype
, t
,
434 fold_convert_loc (loc
, itype
,
436 t
= fold_convert_loc (loc
, long_long_unsigned_type_node
, t
);
437 if (count
!= NULL_TREE
)
438 count
= fold_build2_loc (loc
, MULT_EXPR
,
439 long_long_unsigned_type_node
,
443 if (TREE_CODE (count
) != INTEGER_CST
)
446 else if (count
&& !integer_zerop (count
))
453 && (fd
->sched_kind
!= OMP_CLAUSE_SCHEDULE_STATIC
454 || fd
->have_ordered
))
456 if (!tree_int_cst_lt (count
, TYPE_MAX_VALUE (long_integer_type_node
)))
457 iter_type
= long_long_unsigned_type_node
;
459 iter_type
= long_integer_type_node
;
461 else if (collapse_iter
&& *collapse_iter
!= NULL
)
462 iter_type
= TREE_TYPE (*collapse_iter
);
463 fd
->iter_type
= iter_type
;
464 if (collapse_iter
&& *collapse_iter
== NULL
)
465 *collapse_iter
= create_tmp_var (iter_type
, ".iter");
466 if (collapse_count
&& *collapse_count
== NULL
)
469 *collapse_count
= fold_convert_loc (loc
, iter_type
, count
);
471 *collapse_count
= create_tmp_var (iter_type
, ".count");
474 if (fd
->collapse
> 1 || fd
->tiling
|| (fd
->ordered
&& loops
))
476 fd
->loop
.v
= *collapse_iter
;
477 fd
->loop
.n1
= build_int_cst (TREE_TYPE (fd
->loop
.v
), 0);
478 fd
->loop
.n2
= *collapse_count
;
479 fd
->loop
.step
= build_int_cst (TREE_TYPE (fd
->loop
.v
), 1);
480 fd
->loop
.cond_code
= LT_EXPR
;
486 /* Build a call to GOMP_barrier. */
489 omp_build_barrier (tree lhs
)
491 tree fndecl
= builtin_decl_explicit (lhs
? BUILT_IN_GOMP_BARRIER_CANCEL
492 : BUILT_IN_GOMP_BARRIER
);
493 gcall
*g
= gimple_build_call (fndecl
, 0);
495 gimple_call_set_lhs (g
, lhs
);
499 /* Return maximum possible vectorization factor for the target. */
506 || !flag_tree_loop_optimize
507 || (!flag_tree_loop_vectorize
508 && global_options_set
.x_flag_tree_loop_vectorize
))
511 auto_vector_sizes sizes
;
512 targetm
.vectorize
.autovectorize_vector_sizes (&sizes
, true);
513 if (!sizes
.is_empty ())
516 for (unsigned int i
= 0; i
< sizes
.length (); ++i
)
517 vf
= ordered_max (vf
, sizes
[i
]);
521 machine_mode vqimode
= targetm
.vectorize
.preferred_simd_mode (QImode
);
522 if (GET_MODE_CLASS (vqimode
) == MODE_VECTOR_INT
)
523 return GET_MODE_NUNITS (vqimode
);
528 /* Return maximum SIMT width if offloading may target SIMT hardware. */
531 omp_max_simt_vf (void)
535 if (ENABLE_OFFLOADING
)
536 for (const char *c
= getenv ("OFFLOAD_TARGET_NAMES"); c
;)
538 if (!strncmp (c
, "nvptx", strlen ("nvptx")))
540 else if ((c
= strchr (c
, ',')))
546 /* Store the construct selectors as tree codes from last to first,
547 return their number. */
550 omp_constructor_traits_to_codes (tree ctx
, enum tree_code
*constructs
)
552 int nconstructs
= list_length (ctx
);
553 int i
= nconstructs
- 1;
554 for (tree t2
= ctx
; t2
; t2
= TREE_CHAIN (t2
), i
--)
556 const char *sel
= IDENTIFIER_POINTER (TREE_PURPOSE (t2
));
557 if (!strcmp (sel
, "target"))
558 constructs
[i
] = OMP_TARGET
;
559 else if (!strcmp (sel
, "teams"))
560 constructs
[i
] = OMP_TEAMS
;
561 else if (!strcmp (sel
, "parallel"))
562 constructs
[i
] = OMP_PARALLEL
;
563 else if (!strcmp (sel
, "for") || !strcmp (sel
, "do"))
564 constructs
[i
] = OMP_FOR
;
565 else if (!strcmp (sel
, "simd"))
566 constructs
[i
] = OMP_SIMD
;
570 gcc_assert (i
== -1);
574 /* Return 1 if context selector matches the current OpenMP context, 0
575 if it does not and -1 if it is unknown and need to be determined later.
576 Some properties can be checked right away during parsing (this routine),
577 others need to wait until the whole TU is parsed, others need to wait until
578 IPA, others until vectorization. */
581 omp_context_selector_matches (tree ctx
)
584 for (tree t1
= ctx
; t1
; t1
= TREE_CHAIN (t1
))
586 char set
= IDENTIFIER_POINTER (TREE_PURPOSE (t1
))[0];
589 /* For now, ignore the construct set. While something can be
590 determined already during parsing, we don't know until end of TU
591 whether additional constructs aren't added through declare variant
592 unless "omp declare variant variant" attribute exists already
593 (so in most of the cases), and we'd need to maintain set of
594 surrounding OpenMP constructs, which is better handled during
596 if (symtab
->state
== PARSING
597 || (cfun
->curr_properties
& PROP_gimple_any
) != 0)
603 enum tree_code constructs
[5];
605 = omp_constructor_traits_to_codes (TREE_VALUE (t1
), constructs
);
607 = omp_construct_selector_matches (constructs
, nconstructs
);
614 for (tree t2
= TREE_VALUE (t1
); t2
; t2
= TREE_CHAIN (t2
))
616 const char *sel
= IDENTIFIER_POINTER (TREE_PURPOSE (t2
));
620 if (set
== 'i' && !strcmp (sel
, "vendor"))
621 for (tree t3
= TREE_VALUE (t2
); t3
; t3
= TREE_CHAIN (t3
))
623 const char *prop
= IDENTIFIER_POINTER (TREE_PURPOSE (t3
));
624 if (!strcmp (prop
, " score") || !strcmp (prop
, "gnu"))
630 if (set
== 'i' && !strcmp (sel
, "extension"))
631 /* We don't support any extensions right now. */
635 if (set
== 'i' && !strcmp (sel
, "atomic_default_mem_order"))
637 enum omp_memory_order omo
638 = ((enum omp_memory_order
)
640 & OMP_REQUIRES_ATOMIC_DEFAULT_MEM_ORDER
));
641 if (omo
== OMP_MEMORY_ORDER_UNSPECIFIED
)
643 /* We don't know yet, until end of TU. */
644 if (symtab
->state
== PARSING
)
650 omo
= OMP_MEMORY_ORDER_RELAXED
;
652 tree t3
= TREE_VALUE (t2
);
653 const char *prop
= IDENTIFIER_POINTER (TREE_PURPOSE (t3
));
654 if (!strcmp (prop
, " score"))
656 t3
= TREE_CHAIN (t3
);
657 prop
= IDENTIFIER_POINTER (TREE_PURPOSE (t3
));
659 if (!strcmp (prop
, "relaxed")
660 && omo
!= OMP_MEMORY_ORDER_RELAXED
)
662 else if (!strcmp (prop
, "seq_cst")
663 && omo
!= OMP_MEMORY_ORDER_SEQ_CST
)
665 else if (!strcmp (prop
, "acq_rel")
666 && omo
!= OMP_MEMORY_ORDER_ACQ_REL
)
669 if (set
== 'd' && !strcmp (sel
, "arch"))
670 /* For now, need a target hook. */
674 if (set
== 'i' && !strcmp (sel
, "unified_address"))
676 if ((omp_requires_mask
& OMP_REQUIRES_UNIFIED_ADDRESS
) == 0)
678 if (symtab
->state
== PARSING
)
685 if (set
== 'i' && !strcmp (sel
, "unified_shared_memory"))
687 if ((omp_requires_mask
688 & OMP_REQUIRES_UNIFIED_SHARED_MEMORY
) == 0)
690 if (symtab
->state
== PARSING
)
699 if (set
== 'i' && !strcmp (sel
, "dynamic_allocators"))
701 if ((omp_requires_mask
702 & OMP_REQUIRES_DYNAMIC_ALLOCATORS
) == 0)
704 if (symtab
->state
== PARSING
)
713 if (set
== 'i' && !strcmp (sel
, "reverse_offload"))
715 if ((omp_requires_mask
& OMP_REQUIRES_REVERSE_OFFLOAD
) == 0)
717 if (symtab
->state
== PARSING
)
726 if (set
== 'd' && !strcmp (sel
, "kind"))
727 for (tree t3
= TREE_VALUE (t2
); t3
; t3
= TREE_CHAIN (t3
))
729 const char *prop
= IDENTIFIER_POINTER (TREE_PURPOSE (t3
));
730 if (!strcmp (prop
, "any"))
732 if (!strcmp (prop
, "fpga"))
733 return 0; /* Right now GCC doesn't support any fpgas. */
734 if (!strcmp (prop
, "host"))
736 if (ENABLE_OFFLOADING
|| hsa_gen_requested_p ())
740 if (!strcmp (prop
, "nohost"))
742 if (ENABLE_OFFLOADING
|| hsa_gen_requested_p ())
748 if (!strcmp (prop
, "cpu") || !strcmp (prop
, "gpu"))
750 bool maybe_gpu
= false;
751 if (hsa_gen_requested_p ())
753 else if (ENABLE_OFFLOADING
)
754 for (const char *c
= getenv ("OFFLOAD_TARGET_NAMES");
757 if (!strncmp (c
, "nvptx", strlen ("nvptx"))
758 || !strncmp (c
, "amdgcn", strlen ("amdgcn")))
763 else if ((c
= strchr (c
, ',')))
775 /* Any other kind doesn't match. */
780 if (set
== 'd' && !strcmp (sel
, "isa"))
781 /* For now, need a target hook. */
785 if (set
== 'u' && !strcmp (sel
, "condition"))
786 for (tree t3
= TREE_VALUE (t2
); t3
; t3
= TREE_CHAIN (t3
))
787 if (TREE_PURPOSE (t3
) == NULL_TREE
)
789 if (integer_zerop (TREE_VALUE (t3
)))
791 if (integer_nonzerop (TREE_VALUE (t3
)))
804 /* Try to resolve declare variant, return the variant decl if it should
805 be used instead of base, or base otherwise. */
808 omp_resolve_declare_variant (tree base
)
810 tree variant
= NULL_TREE
;
811 for (tree attr
= DECL_ATTRIBUTES (base
); attr
; attr
= TREE_CHAIN (attr
))
813 attr
= lookup_attribute ("omp declare variant base", attr
);
814 if (attr
== NULL_TREE
)
816 switch (omp_context_selector_matches (TREE_VALUE (TREE_VALUE (attr
))))
819 /* No match, ignore. */
822 /* Needs to be deferred. */
825 /* FIXME: Scoring not implemented yet, so just resolve it
826 if there is a single variant only. */
829 if (TREE_CODE (TREE_PURPOSE (TREE_VALUE (attr
))) == FUNCTION_DECL
)
830 variant
= TREE_PURPOSE (TREE_VALUE (attr
));
835 return variant
? variant
: base
;
839 /* Encode an oacc launch argument. This matches the GOMP_LAUNCH_PACK
840 macro on gomp-constants.h. We do not check for overflow. */
843 oacc_launch_pack (unsigned code
, tree device
, unsigned op
)
847 res
= build_int_cst (unsigned_type_node
, GOMP_LAUNCH_PACK (code
, 0, op
));
850 device
= fold_build2 (LSHIFT_EXPR
, unsigned_type_node
,
851 device
, build_int_cst (unsigned_type_node
,
852 GOMP_LAUNCH_DEVICE_SHIFT
));
853 res
= fold_build2 (BIT_IOR_EXPR
, unsigned_type_node
, res
, device
);
858 /* FIXME: What is the following comment for? */
859 /* Look for compute grid dimension clauses and convert to an attribute
860 attached to FN. This permits the target-side code to (a) massage
861 the dimensions, (b) emit that data and (c) optimize. Non-constant
862 dimensions are pushed onto ARGS.
864 The attribute value is a TREE_LIST. A set of dimensions is
865 represented as a list of INTEGER_CST. Those that are runtime
866 exprs are represented as an INTEGER_CST of zero.
868 TODO: Normally the attribute will just contain a single such list. If
869 however it contains a list of lists, this will represent the use of
870 device_type. Each member of the outer list is an assoc list of
871 dimensions, keyed by the device type. The first entry will be the
872 default. Well, that's the plan. */
874 /* Replace any existing oacc fn attribute with updated dimensions. */
876 /* Variant working on a list of attributes. */
879 oacc_replace_fn_attrib_attr (tree attribs
, tree dims
)
881 tree ident
= get_identifier (OACC_FN_ATTRIB
);
883 /* If we happen to be present as the first attrib, drop it. */
884 if (attribs
&& TREE_PURPOSE (attribs
) == ident
)
885 attribs
= TREE_CHAIN (attribs
);
886 return tree_cons (ident
, dims
, attribs
);
889 /* Variant working on a function decl. */
892 oacc_replace_fn_attrib (tree fn
, tree dims
)
895 = oacc_replace_fn_attrib_attr (DECL_ATTRIBUTES (fn
), dims
);
898 /* Scan CLAUSES for launch dimensions and attach them to the oacc
899 function attribute. Push any that are non-constant onto the ARGS
900 list, along with an appropriate GOMP_LAUNCH_DIM tag. */
903 oacc_set_fn_attrib (tree fn
, tree clauses
, vec
<tree
> *args
)
905 /* Must match GOMP_DIM ordering. */
906 static const omp_clause_code ids
[]
907 = { OMP_CLAUSE_NUM_GANGS
, OMP_CLAUSE_NUM_WORKERS
,
908 OMP_CLAUSE_VECTOR_LENGTH
};
910 tree dims
[GOMP_DIM_MAX
];
912 tree attr
= NULL_TREE
;
913 unsigned non_const
= 0;
915 for (ix
= GOMP_DIM_MAX
; ix
--;)
917 tree clause
= omp_find_clause (clauses
, ids
[ix
]);
918 tree dim
= NULL_TREE
;
921 dim
= OMP_CLAUSE_EXPR (clause
, ids
[ix
]);
923 if (dim
&& TREE_CODE (dim
) != INTEGER_CST
)
925 dim
= integer_zero_node
;
926 non_const
|= GOMP_DIM_MASK (ix
);
928 attr
= tree_cons (NULL_TREE
, dim
, attr
);
931 oacc_replace_fn_attrib (fn
, attr
);
935 /* Push a dynamic argument set. */
936 args
->safe_push (oacc_launch_pack (GOMP_LAUNCH_DIM
,
937 NULL_TREE
, non_const
));
938 for (unsigned ix
= 0; ix
!= GOMP_DIM_MAX
; ix
++)
939 if (non_const
& GOMP_DIM_MASK (ix
))
940 args
->safe_push (dims
[ix
]);
944 /* Verify OpenACC routine clauses.
946 Returns 0 if FNDECL should be marked with an OpenACC 'routine' directive, 1
947 if it has already been marked in compatible way, and -1 if incompatible.
948 Upon returning, the chain of clauses will contain exactly one clause
949 specifying the level of parallelism. */
952 oacc_verify_routine_clauses (tree fndecl
, tree
*clauses
, location_t loc
,
953 const char *routine_str
)
955 tree c_level
= NULL_TREE
;
956 tree c_p
= NULL_TREE
;
957 for (tree c
= *clauses
; c
; c_p
= c
, c
= OMP_CLAUSE_CHAIN (c
))
958 switch (OMP_CLAUSE_CODE (c
))
960 case OMP_CLAUSE_GANG
:
961 case OMP_CLAUSE_WORKER
:
962 case OMP_CLAUSE_VECTOR
:
964 if (c_level
== NULL_TREE
)
966 else if (OMP_CLAUSE_CODE (c
) == OMP_CLAUSE_CODE (c_level
))
968 /* This has already been diagnosed in the front ends. */
969 /* Drop the duplicate clause. */
970 gcc_checking_assert (c_p
!= NULL_TREE
);
971 OMP_CLAUSE_CHAIN (c_p
) = OMP_CLAUSE_CHAIN (c
);
976 error_at (OMP_CLAUSE_LOCATION (c
),
977 "%qs specifies a conflicting level of parallelism",
978 omp_clause_code_name
[OMP_CLAUSE_CODE (c
)]);
979 inform (OMP_CLAUSE_LOCATION (c_level
),
980 "... to the previous %qs clause here",
981 omp_clause_code_name
[OMP_CLAUSE_CODE (c_level
)]);
982 /* Drop the conflicting clause. */
983 gcc_checking_assert (c_p
!= NULL_TREE
);
984 OMP_CLAUSE_CHAIN (c_p
) = OMP_CLAUSE_CHAIN (c
);
991 if (c_level
== NULL_TREE
)
993 /* Default to an implicit 'seq' clause. */
994 c_level
= build_omp_clause (loc
, OMP_CLAUSE_SEQ
);
995 OMP_CLAUSE_CHAIN (c_level
) = *clauses
;
998 /* In *clauses, we now have exactly one clause specifying the level of
1002 = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (fndecl
));
1003 if (attr
!= NULL_TREE
)
1005 /* If a "#pragma acc routine" has already been applied, just verify
1006 this one for compatibility. */
1007 /* Collect previous directive's clauses. */
1008 tree c_level_p
= NULL_TREE
;
1009 for (tree c
= TREE_VALUE (attr
); c
; c
= OMP_CLAUSE_CHAIN (c
))
1010 switch (OMP_CLAUSE_CODE (c
))
1012 case OMP_CLAUSE_GANG
:
1013 case OMP_CLAUSE_WORKER
:
1014 case OMP_CLAUSE_VECTOR
:
1015 case OMP_CLAUSE_SEQ
:
1016 gcc_checking_assert (c_level_p
== NULL_TREE
);
1022 gcc_checking_assert (c_level_p
!= NULL_TREE
);
1023 /* ..., and compare to current directive's, which we've already collected
1027 /* Matching level of parallelism? */
1028 if (OMP_CLAUSE_CODE (c_level
) != OMP_CLAUSE_CODE (c_level_p
))
1031 c_diag_p
= c_level_p
;
1038 if (c_diag
!= NULL_TREE
)
1039 error_at (OMP_CLAUSE_LOCATION (c_diag
),
1040 "incompatible %qs clause when applying"
1041 " %<%s%> to %qD, which has already been"
1042 " marked with an OpenACC 'routine' directive",
1043 omp_clause_code_name
[OMP_CLAUSE_CODE (c_diag
)],
1044 routine_str
, fndecl
);
1045 else if (c_diag_p
!= NULL_TREE
)
1047 "missing %qs clause when applying"
1048 " %<%s%> to %qD, which has already been"
1049 " marked with an OpenACC 'routine' directive",
1050 omp_clause_code_name
[OMP_CLAUSE_CODE (c_diag_p
)],
1051 routine_str
, fndecl
);
1054 if (c_diag_p
!= NULL_TREE
)
1055 inform (OMP_CLAUSE_LOCATION (c_diag_p
),
1056 "... with %qs clause here",
1057 omp_clause_code_name
[OMP_CLAUSE_CODE (c_diag_p
)]);
1060 /* In the front ends, we don't preserve location information for the
1061 OpenACC routine directive itself. However, that of c_level_p
1063 location_t loc_routine
= OMP_CLAUSE_LOCATION (c_level_p
);
1064 inform (loc_routine
, "... without %qs clause near to here",
1065 omp_clause_code_name
[OMP_CLAUSE_CODE (c_diag
)]);
1074 /* Process the OpenACC 'routine' directive clauses to generate an attribute
1075 for the level of parallelism. All dimensions have a size of zero
1076 (dynamic). TREE_PURPOSE is set to indicate whether that dimension
1077 can have a loop partitioned on it. non-zero indicates
1078 yes, zero indicates no. By construction once a non-zero has been
1079 reached, further inner dimensions must also be non-zero. We set
1080 TREE_VALUE to zero for the dimensions that may be partitioned and
1081 1 for the other ones -- if a loop is (erroneously) spawned at
1082 an outer level, we don't want to try and partition it. */
1085 oacc_build_routine_dims (tree clauses
)
1087 /* Must match GOMP_DIM ordering. */
1088 static const omp_clause_code ids
[]
1089 = {OMP_CLAUSE_GANG
, OMP_CLAUSE_WORKER
, OMP_CLAUSE_VECTOR
, OMP_CLAUSE_SEQ
};
1093 for (; clauses
; clauses
= OMP_CLAUSE_CHAIN (clauses
))
1094 for (ix
= GOMP_DIM_MAX
+ 1; ix
--;)
1095 if (OMP_CLAUSE_CODE (clauses
) == ids
[ix
])
1100 gcc_checking_assert (level
>= 0);
1102 tree dims
= NULL_TREE
;
1104 for (ix
= GOMP_DIM_MAX
; ix
--;)
1105 dims
= tree_cons (build_int_cst (boolean_type_node
, ix
>= level
),
1106 build_int_cst (integer_type_node
, ix
< level
), dims
);
1111 /* Retrieve the oacc function attrib and return it. Non-oacc
1112 functions will return NULL. */
1115 oacc_get_fn_attrib (tree fn
)
1117 return lookup_attribute (OACC_FN_ATTRIB
, DECL_ATTRIBUTES (fn
));
1120 /* Return true if FN is an OpenMP or OpenACC offloading function. */
1123 offloading_function_p (tree fn
)
1125 tree attrs
= DECL_ATTRIBUTES (fn
);
1126 return (lookup_attribute ("omp declare target", attrs
)
1127 || lookup_attribute ("omp target entrypoint", attrs
));
1130 /* Extract an oacc execution dimension from FN. FN must be an
1131 offloaded function or routine that has already had its execution
1132 dimensions lowered to the target-specific values. */
1135 oacc_get_fn_dim_size (tree fn
, int axis
)
1137 tree attrs
= oacc_get_fn_attrib (fn
);
1139 gcc_assert (axis
< GOMP_DIM_MAX
);
1141 tree dims
= TREE_VALUE (attrs
);
1143 dims
= TREE_CHAIN (dims
);
1145 int size
= TREE_INT_CST_LOW (TREE_VALUE (dims
));
1150 /* Extract the dimension axis from an IFN_GOACC_DIM_POS or
1151 IFN_GOACC_DIM_SIZE call. */
1154 oacc_get_ifn_dim_arg (const gimple
*stmt
)
1156 gcc_checking_assert (gimple_call_internal_fn (stmt
) == IFN_GOACC_DIM_SIZE
1157 || gimple_call_internal_fn (stmt
) == IFN_GOACC_DIM_POS
);
1158 tree arg
= gimple_call_arg (stmt
, 0);
1159 HOST_WIDE_INT axis
= TREE_INT_CST_LOW (arg
);
1161 gcc_checking_assert (axis
>= 0 && axis
< GOMP_DIM_MAX
);