?  search:
summary | log | graphiclog1 | graphiclog2 | commit | commitdiff | tree | refs | edit | fork
blame | history | raw | HEAD
fix __builtin___clear_cache overrider fallout
[official-gcc.git] / gcc / omp-general.c
blobb66dfb58257812139687334aafd18d89b29345d3
1 /* General types and functions that are uselful for processing of OpenMP,
2 OpenACC and similar directivers at various stages of compilation.
3
4 Copyright (C) 2005-2020 Free Software Foundation, Inc.
5
6 This file is part of GCC.
7
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
11 version.
12
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
16 for more details.
17
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/>. */
21
22 /* Find an OMP clause of type KIND within CLAUSES. */
23
24 #include "config.h"
25 #include "system.h"
26 #include "coretypes.h"
27 #include "backend.h"
28 #include "target.h"
29 #include "tree.h"
30 #include "gimple.h"
31 #include "ssa.h"
32 #include "diagnostic-core.h"
33 #include "fold-const.h"
34 #include "langhooks.h"
35 #include "omp-general.h"
36 #include "stringpool.h"
37 #include "attribs.h"
38 #include "gimplify.h"
39 #include "cgraph.h"
40 #include "alloc-pool.h"
41 #include "symbol-summary.h"
42 #include "tree-pass.h"
43 #include "omp-device-properties.h"
44 #include "tree-iterator.h"
45 #include "data-streamer.h"
46 #include "streamer-hooks.h"
47
48 enum omp_requires omp_requires_mask;
49
50 tree
51 omp_find_clause (tree clauses, enum omp_clause_code kind)
52 {
53 for (; clauses ; clauses = OMP_CLAUSE_CHAIN (clauses))
54 if (OMP_CLAUSE_CODE (clauses) == kind)
55 return clauses;
56
57 return NULL_TREE;
58 }
59
60 /* True if OpenMP should regard this DECL as being a scalar which has Fortran's
61 allocatable or pointer attribute. */
62 bool
63 omp_is_allocatable_or_ptr (tree decl)
64 {
65 return lang_hooks.decls.omp_is_allocatable_or_ptr (decl);
66 }
67
68 /* Check whether this DECL belongs to a Fortran optional argument.
69 With 'for_present_check' set to false, decls which are optional parameters
70 themselve are returned as tree - or a NULL_TREE otherwise. Those decls are
71 always pointers. With 'for_present_check' set to true, the decl for checking
72 whether an argument is present is returned; for arguments with value
73 attribute this is the hidden argument and of BOOLEAN_TYPE. If the decl is
74 unrelated to optional arguments, NULL_TREE is returned. */
75
76 tree
77 omp_check_optional_argument (tree decl, bool for_present_check)
78 {
79 return lang_hooks.decls.omp_check_optional_argument (decl, for_present_check);
80 }
81
82 /* Return true if DECL is a reference type. */
83
84 bool
85 omp_is_reference (tree decl)
86 {
87 return lang_hooks.decls.omp_privatize_by_reference (decl);
88 }
89
90 /* Adjust *COND_CODE and *N2 so that the former is either LT_EXPR or GT_EXPR,
91 given that V is the loop index variable and STEP is loop step. */
92
93 void
94 omp_adjust_for_condition (location_t loc, enum tree_code *cond_code, tree *n2,
95 tree v, tree step)
96 {
97 switch (*cond_code)
98 {
99 case LT_EXPR:
100 case GT_EXPR:
101 break;
102
103 case NE_EXPR:
104 gcc_assert (TREE_CODE (step) == INTEGER_CST);
105 if (TREE_CODE (TREE_TYPE (v)) == INTEGER_TYPE)
106 {
107 if (integer_onep (step))
108 *cond_code = LT_EXPR;
109 else
110 {
111 gcc_assert (integer_minus_onep (step));
112 *cond_code = GT_EXPR;
113 }
114 }
115 else
116 {
117 tree unit = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (v)));
118 gcc_assert (TREE_CODE (unit) == INTEGER_CST);
119 if (tree_int_cst_equal (unit, step))
120 *cond_code = LT_EXPR;
121 else
122 {
123 gcc_assert (wi::neg (wi::to_widest (unit))
124 == wi::to_widest (step));
125 *cond_code = GT_EXPR;
126 }
127 }
128
129 break;
130
131 case LE_EXPR:
132 if (POINTER_TYPE_P (TREE_TYPE (*n2)))
133 *n2 = fold_build_pointer_plus_hwi_loc (loc, *n2, 1);
134 else
135 *n2 = fold_build2_loc (loc, PLUS_EXPR, TREE_TYPE (*n2), *n2,
136 build_int_cst (TREE_TYPE (*n2), 1));
137 *cond_code = LT_EXPR;
138 break;
139 case GE_EXPR:
140 if (POINTER_TYPE_P (TREE_TYPE (*n2)))
141 *n2 = fold_build_pointer_plus_hwi_loc (loc, *n2, -1);
142 else
143 *n2 = fold_build2_loc (loc, MINUS_EXPR, TREE_TYPE (*n2), *n2,
144 build_int_cst (TREE_TYPE (*n2), 1));
145 *cond_code = GT_EXPR;
146 break;
147 default:
148 gcc_unreachable ();
149 }
150 }
151
152 /* Return the looping step from INCR, extracted from the step of a gimple omp
153 for statement. */
154
155 tree
156 omp_get_for_step_from_incr (location_t loc, tree incr)
157 {
158 tree step;
159 switch (TREE_CODE (incr))
160 {
161 case PLUS_EXPR:
162 step = TREE_OPERAND (incr, 1);
163 break;
164 case POINTER_PLUS_EXPR:
165 step = fold_convert (ssizetype, TREE_OPERAND (incr, 1));
166 break;
167 case MINUS_EXPR:
168 step = TREE_OPERAND (incr, 1);
169 step = fold_build1_loc (loc, NEGATE_EXPR, TREE_TYPE (step), step);
170 break;
171 default:
172 gcc_unreachable ();
173 }
174 return step;
175 }
176
177 /* Extract the header elements of parallel loop FOR_STMT and store
178 them into *FD. */
179
180 void
181 omp_extract_for_data (gomp_for *for_stmt, struct omp_for_data *fd,
182 struct omp_for_data_loop *loops)
183 {
184 tree t, var, *collapse_iter, *collapse_count;
185 tree count = NULL_TREE, iter_type = long_integer_type_node;
186 struct omp_for_data_loop *loop;
187 int i;
188 struct omp_for_data_loop dummy_loop;
189 location_t loc = gimple_location (for_stmt);
190 bool simd = gimple_omp_for_kind (for_stmt) == GF_OMP_FOR_KIND_SIMD;
191 bool distribute = gimple_omp_for_kind (for_stmt)
192 == GF_OMP_FOR_KIND_DISTRIBUTE;
193 bool taskloop = gimple_omp_for_kind (for_stmt)
194 == GF_OMP_FOR_KIND_TASKLOOP;
195 tree iterv, countv;
196
197 fd->for_stmt = for_stmt;
198 fd->pre = NULL;
199 fd->have_nowait = distribute || simd;
200 fd->have_ordered = false;
201 fd->have_reductemp = false;
202 fd->have_pointer_condtemp = false;
203 fd->have_scantemp = false;
204 fd->have_nonctrl_scantemp = false;
205 fd->non_rect = false;
206 fd->lastprivate_conditional = 0;
207 fd->tiling = NULL_TREE;
208 fd->collapse = 1;
209 fd->ordered = 0;
210 fd->first_nonrect = -1;
211 fd->last_nonrect = -1;
212 fd->sched_kind = OMP_CLAUSE_SCHEDULE_STATIC;
213 fd->sched_modifiers = 0;
214 fd->chunk_size = NULL_TREE;
215 fd->simd_schedule = false;
216 fd->first_inner_iterations = NULL_TREE;
217 fd->factor = NULL_TREE;
218 fd->adjn1 = NULL_TREE;
219 collapse_iter = NULL;
220 collapse_count = NULL;
221
222 for (t = gimple_omp_for_clauses (for_stmt); t ; t = OMP_CLAUSE_CHAIN (t))
223 switch (OMP_CLAUSE_CODE (t))
224 {
225 case OMP_CLAUSE_NOWAIT:
226 fd->have_nowait = true;
227 break;
228 case OMP_CLAUSE_ORDERED:
229 fd->have_ordered = true;
230 if (OMP_CLAUSE_ORDERED_EXPR (t))
231 fd->ordered = tree_to_shwi (OMP_CLAUSE_ORDERED_EXPR (t));
232 break;
233 case OMP_CLAUSE_SCHEDULE:
234 gcc_assert (!distribute && !taskloop);
235 fd->sched_kind
236 = (enum omp_clause_schedule_kind)
237 (OMP_CLAUSE_SCHEDULE_KIND (t) & OMP_CLAUSE_SCHEDULE_MASK);
238 fd->sched_modifiers = (OMP_CLAUSE_SCHEDULE_KIND (t)
239 & ~OMP_CLAUSE_SCHEDULE_MASK);
240 fd->chunk_size = OMP_CLAUSE_SCHEDULE_CHUNK_EXPR (t);
241 fd->simd_schedule = OMP_CLAUSE_SCHEDULE_SIMD (t);
242 break;
243 case OMP_CLAUSE_DIST_SCHEDULE:
244 gcc_assert (distribute);
245 fd->chunk_size = OMP_CLAUSE_DIST_SCHEDULE_CHUNK_EXPR (t);
246 break;
247 case OMP_CLAUSE_COLLAPSE:
248 fd->collapse = tree_to_shwi (OMP_CLAUSE_COLLAPSE_EXPR (t));
249 if (fd->collapse > 1)
250 {
251 collapse_iter = &OMP_CLAUSE_COLLAPSE_ITERVAR (t);
252 collapse_count = &OMP_CLAUSE_COLLAPSE_COUNT (t);
253 }
254 break;
255 case OMP_CLAUSE_TILE:
256 fd->tiling = OMP_CLAUSE_TILE_LIST (t);
257 fd->collapse = list_length (fd->tiling);
258 gcc_assert (fd->collapse);
259 collapse_iter = &OMP_CLAUSE_TILE_ITERVAR (t);
260 collapse_count = &OMP_CLAUSE_TILE_COUNT (t);
261 break;
262 case OMP_CLAUSE__REDUCTEMP_:
263 fd->have_reductemp = true;
264 break;
265 case OMP_CLAUSE_LASTPRIVATE:
266 if (OMP_CLAUSE_LASTPRIVATE_CONDITIONAL (t))
267 fd->lastprivate_conditional++;
268 break;
269 case OMP_CLAUSE__CONDTEMP_:
270 if (POINTER_TYPE_P (TREE_TYPE (OMP_CLAUSE_DECL (t))))
271 fd->have_pointer_condtemp = true;
272 break;
273 case OMP_CLAUSE__SCANTEMP_:
274 fd->have_scantemp = true;
275 if (!OMP_CLAUSE__SCANTEMP__ALLOC (t)
276 && !OMP_CLAUSE__SCANTEMP__CONTROL (t))
277 fd->have_nonctrl_scantemp = true;
278 break;
279 default:
280 break;
281 }
282
283 if (fd->collapse > 1 || fd->tiling)
284 fd->loops = loops;
285 else
286 fd->loops = &fd->loop;
287
288 if (fd->ordered && fd->collapse == 1 && loops != NULL)
289 {
290 fd->loops = loops;
291 iterv = NULL_TREE;
292 countv = NULL_TREE;
293 collapse_iter = &iterv;
294 collapse_count = &countv;
295 }
296
297 /* FIXME: for now map schedule(auto) to schedule(static).
298 There should be analysis to determine whether all iterations
299 are approximately the same amount of work (then schedule(static)
300 is best) or if it varies (then schedule(dynamic,N) is better). */
301 if (fd->sched_kind == OMP_CLAUSE_SCHEDULE_AUTO)
302 {
303 fd->sched_kind = OMP_CLAUSE_SCHEDULE_STATIC;
304 gcc_assert (fd->chunk_size == NULL);
305 }
306 gcc_assert ((fd->collapse == 1 && !fd->tiling) || collapse_iter != NULL);
307 if (taskloop)
308 fd->sched_kind = OMP_CLAUSE_SCHEDULE_RUNTIME;
309 if (fd->sched_kind == OMP_CLAUSE_SCHEDULE_RUNTIME)
310 gcc_assert (fd->chunk_size == NULL);
311 else if (fd->chunk_size == NULL)
312 {
313 /* We only need to compute a default chunk size for ordered
314 static loops and dynamic loops. */
315 if (fd->sched_kind != OMP_CLAUSE_SCHEDULE_STATIC
316 || fd->have_ordered)
317 fd->chunk_size = (fd->sched_kind == OMP_CLAUSE_SCHEDULE_STATIC)
318 ? integer_zero_node : integer_one_node;
319 }
320
321 int cnt = fd->ordered ? fd->ordered : fd->collapse;
322 int single_nonrect = -1;
323 tree single_nonrect_count = NULL_TREE;
324 enum tree_code single_nonrect_cond_code = ERROR_MARK;
325 for (i = 1; i < cnt; i++)
326 {
327 tree n1 = gimple_omp_for_initial (for_stmt, i);
328 tree n2 = gimple_omp_for_final (for_stmt, i);
329 if (TREE_CODE (n1) == TREE_VEC)
330 {
331 if (fd->non_rect)
332 {
333 single_nonrect = -1;
334 break;
335 }
336 for (int j = i - 1; j >= 0; j--)
337 if (TREE_VEC_ELT (n1, 0) == gimple_omp_for_index (for_stmt, j))
338 {
339 single_nonrect = j;
340 break;
341 }
342 fd->non_rect = true;
343 }
344 else if (TREE_CODE (n2) == TREE_VEC)
345 {
346 if (fd->non_rect)
347 {
348 single_nonrect = -1;
349 break;
350 }
351 for (int j = i - 1; j >= 0; j--)
352 if (TREE_VEC_ELT (n2, 0) == gimple_omp_for_index (for_stmt, j))
353 {
354 single_nonrect = j;
355 break;
356 }
357 fd->non_rect = true;
358 }
359 }
360 for (i = 0; i < cnt; i++)
361 {
362 if (i == 0
363 && fd->collapse == 1
364 && !fd->tiling
365 && (fd->ordered == 0 || loops == NULL))
366 loop = &fd->loop;
367 else if (loops != NULL)
368 loop = loops + i;
369 else
370 loop = &dummy_loop;
371
372 loop->v = gimple_omp_for_index (for_stmt, i);
373 gcc_assert (SSA_VAR_P (loop->v));
374 gcc_assert (TREE_CODE (TREE_TYPE (loop->v)) == INTEGER_TYPE
375 || TREE_CODE (TREE_TYPE (loop->v)) == POINTER_TYPE);
376 var = TREE_CODE (loop->v) == SSA_NAME ? SSA_NAME_VAR (loop->v) : loop->v;
377 loop->n1 = gimple_omp_for_initial (for_stmt, i);
378 loop->m1 = NULL_TREE;
379 loop->m2 = NULL_TREE;
380 loop->outer = 0;
381 loop->non_rect_referenced = false;
382 if (TREE_CODE (loop->n1) == TREE_VEC)
383 {
384 for (int j = i - 1; j >= 0; j--)
385 if (TREE_VEC_ELT (loop->n1, 0) == gimple_omp_for_index (for_stmt, j))
386 {
387 loop->outer = i - j;
388 if (loops != NULL)
389 loops[j].non_rect_referenced = true;
390 if (fd->first_nonrect == -1 || fd->first_nonrect > j)
391 fd->first_nonrect = j;
392 break;
393 }
394 gcc_assert (loop->outer);
395 loop->m1 = TREE_VEC_ELT (loop->n1, 1);
396 loop->n1 = TREE_VEC_ELT (loop->n1, 2);
397 fd->non_rect = true;
398 fd->last_nonrect = i;
399 }
400
401 loop->cond_code = gimple_omp_for_cond (for_stmt, i);
402 loop->n2 = gimple_omp_for_final (for_stmt, i);
403 gcc_assert (loop->cond_code != NE_EXPR
404 || (gimple_omp_for_kind (for_stmt)
405 != GF_OMP_FOR_KIND_OACC_LOOP));
406 if (TREE_CODE (loop->n2) == TREE_VEC)
407 {
408 if (loop->outer)
409 gcc_assert (TREE_VEC_ELT (loop->n2, 0)
410 == gimple_omp_for_index (for_stmt, i - loop->outer));
411 else
412 for (int j = i - 1; j >= 0; j--)
413 if (TREE_VEC_ELT (loop->n2, 0) == gimple_omp_for_index (for_stmt, j))
414 {
415 loop->outer = i - j;
416 if (loops != NULL)
417 loops[j].non_rect_referenced = true;
418 if (fd->first_nonrect == -1 || fd->first_nonrect > j)
419 fd->first_nonrect = j;
420 break;
421 }
422 gcc_assert (loop->outer);
423 loop->m2 = TREE_VEC_ELT (loop->n2, 1);
424 loop->n2 = TREE_VEC_ELT (loop->n2, 2);
425 fd->non_rect = true;
426 fd->last_nonrect = i;
427 }
428
429 t = gimple_omp_for_incr (for_stmt, i);
430 gcc_assert (TREE_OPERAND (t, 0) == var);
431 loop->step = omp_get_for_step_from_incr (loc, t);
432
433 omp_adjust_for_condition (loc, &loop->cond_code, &loop->n2, loop->v,
434 loop->step);
435
436 if (simd
437 || (fd->sched_kind == OMP_CLAUSE_SCHEDULE_STATIC
438 && !fd->have_ordered))
439 {
440 if (fd->collapse == 1 && !fd->tiling)
441 iter_type = TREE_TYPE (loop->v);
442 else if (i == 0
443 || TYPE_PRECISION (iter_type)
444 < TYPE_PRECISION (TREE_TYPE (loop->v)))
445 iter_type
446 = build_nonstandard_integer_type
447 (TYPE_PRECISION (TREE_TYPE (loop->v)), 1);
448 }
449 else if (iter_type != long_long_unsigned_type_node)
450 {
451 if (POINTER_TYPE_P (TREE_TYPE (loop->v)))
452 iter_type = long_long_unsigned_type_node;
453 else if (TYPE_UNSIGNED (TREE_TYPE (loop->v))
454 && TYPE_PRECISION (TREE_TYPE (loop->v))
455 >= TYPE_PRECISION (iter_type))
456 {
457 tree n;
458
459 if (loop->cond_code == LT_EXPR)
460 n = fold_build2_loc (loc, PLUS_EXPR, TREE_TYPE (loop->v),
461 loop->n2, loop->step);
462 else
463 n = loop->n1;
464 if (loop->m1
465 || loop->m2
466 || TREE_CODE (n) != INTEGER_CST
467 || tree_int_cst_lt (TYPE_MAX_VALUE (iter_type), n))
468 iter_type = long_long_unsigned_type_node;
469 }
470 else if (TYPE_PRECISION (TREE_TYPE (loop->v))
471 > TYPE_PRECISION (iter_type))
472 {
473 tree n1, n2;
474
475 if (loop->cond_code == LT_EXPR)
476 {
477 n1 = loop->n1;
478 n2 = fold_build2_loc (loc, PLUS_EXPR, TREE_TYPE (loop->v),
479 loop->n2, loop->step);
480 }
481 else
482 {
483 n1 = fold_build2_loc (loc, MINUS_EXPR, TREE_TYPE (loop->v),
484 loop->n2, loop->step);
485 n2 = loop->n1;
486 }
487 if (loop->m1
488 || loop->m2
489 || TREE_CODE (n1) != INTEGER_CST
490 || TREE_CODE (n2) != INTEGER_CST
491 || !tree_int_cst_lt (TYPE_MIN_VALUE (iter_type), n1)
492 || !tree_int_cst_lt (n2, TYPE_MAX_VALUE (iter_type)))
493 iter_type = long_long_unsigned_type_node;
494 }
495 }
496
497 if (i >= fd->collapse)
498 continue;
499
500 if (collapse_count && *collapse_count == NULL)
501 {
502 if (count && integer_zerop (count))
503 continue;
504 tree n1first = NULL_TREE, n2first = NULL_TREE;
505 tree n1last = NULL_TREE, n2last = NULL_TREE;
506 tree ostep = NULL_TREE;
507 if (loop->m1 || loop->m2)
508 {
509 if (count == NULL_TREE)
510 continue;
511 if (single_nonrect == -1
512 || (loop->m1 && TREE_CODE (loop->m1) != INTEGER_CST)
513 || (loop->m2 && TREE_CODE (loop->m2) != INTEGER_CST)
514 || TREE_CODE (loop->n1) != INTEGER_CST
515 || TREE_CODE (loop->n2) != INTEGER_CST
516 || TREE_CODE (loop->step) != INTEGER_CST)
517 {
518 count = NULL_TREE;
519 continue;
520 }
521 tree var = gimple_omp_for_initial (for_stmt, single_nonrect);
522 tree itype = TREE_TYPE (var);
523 tree first = gimple_omp_for_initial (for_stmt, single_nonrect);
524 t = gimple_omp_for_incr (for_stmt, single_nonrect);
525 ostep = omp_get_for_step_from_incr (loc, t);
526 t = fold_binary (MINUS_EXPR, long_long_unsigned_type_node,
527 single_nonrect_count,
528 build_one_cst (long_long_unsigned_type_node));
529 t = fold_convert (itype, t);
530 first = fold_convert (itype, first);
531 ostep = fold_convert (itype, ostep);
532 tree last = fold_binary (PLUS_EXPR, itype, first,
533 fold_binary (MULT_EXPR, itype, t,
534 ostep));
535 if (TREE_CODE (first) != INTEGER_CST
536 || TREE_CODE (last) != INTEGER_CST)
537 {
538 count = NULL_TREE;
539 continue;
540 }
541 if (loop->m1)
542 {
543 tree m1 = fold_convert (itype, loop->m1);
544 tree n1 = fold_convert (itype, loop->n1);
545 n1first = fold_binary (PLUS_EXPR, itype,
546 fold_binary (MULT_EXPR, itype,
547 first, m1), n1);
548 n1last = fold_binary (PLUS_EXPR, itype,
549 fold_binary (MULT_EXPR, itype,
550 last, m1), n1);
551 }
552 else
553 n1first = n1last = loop->n1;
554 if (loop->m2)
555 {
556 tree n2 = fold_convert (itype, loop->n2);
557 tree m2 = fold_convert (itype, loop->m2);
558 n2first = fold_binary (PLUS_EXPR, itype,
559 fold_binary (MULT_EXPR, itype,
560 first, m2), n2);
561 n2last = fold_binary (PLUS_EXPR, itype,
562 fold_binary (MULT_EXPR, itype,
563 last, m2), n2);
564 }
565 else
566 n2first = n2last = loop->n2;
567 n1first = fold_convert (TREE_TYPE (loop->v), n1first);
568 n2first = fold_convert (TREE_TYPE (loop->v), n2first);
569 n1last = fold_convert (TREE_TYPE (loop->v), n1last);
570 n2last = fold_convert (TREE_TYPE (loop->v), n2last);
571 t = fold_binary (loop->cond_code, boolean_type_node,
572 n1first, n2first);
573 tree t2 = fold_binary (loop->cond_code, boolean_type_node,
574 n1last, n2last);
575 if (t && t2 && integer_nonzerop (t) && integer_nonzerop (t2))
576 /* All outer loop iterators have at least one inner loop
577 iteration. Try to compute the count at compile time. */
578 t = NULL_TREE;
579 else if (t && t2 && integer_zerop (t) && integer_zerop (t2))
580 /* No iterations of the inner loop. count will be set to
581 zero cst below. */;
582 else if (TYPE_UNSIGNED (itype)
583 || t == NULL_TREE
584 || t2 == NULL_TREE
585 || TREE_CODE (t) != INTEGER_CST
586 || TREE_CODE (t2) != INTEGER_CST)
587 {
588 /* Punt (for now). */
589 count = NULL_TREE;
590 continue;
591 }
592 else
593 {
594 /* Some iterations of the outer loop have zero iterations
595 of the inner loop, while others have at least one.
596 In this case, we need to adjust one of those outer
597 loop bounds. If ADJ_FIRST, we need to adjust outer n1
598 (first), otherwise outer n2 (last). */
599 bool adj_first = integer_zerop (t);
600 tree n1 = fold_convert (itype, loop->n1);
601 tree n2 = fold_convert (itype, loop->n2);
602 tree m1 = loop->m1 ? fold_convert (itype, loop->m1)
603 : build_zero_cst (itype);
604 tree m2 = loop->m2 ? fold_convert (itype, loop->m2)
605 : build_zero_cst (itype);
606 t = fold_binary (MINUS_EXPR, itype, n1, n2);
607 t2 = fold_binary (MINUS_EXPR, itype, m2, m1);
608 t = fold_binary (TRUNC_DIV_EXPR, itype, t, t2);
609 t2 = fold_binary (MINUS_EXPR, itype, t, first);
610 t2 = fold_binary (TRUNC_MOD_EXPR, itype, t2, ostep);
611 t = fold_binary (MINUS_EXPR, itype, t, t2);
612 tree n1cur
613 = fold_binary (PLUS_EXPR, itype, n1,
614 fold_binary (MULT_EXPR, itype, m1, t));
615 tree n2cur
616 = fold_binary (PLUS_EXPR, itype, n2,
617 fold_binary (MULT_EXPR, itype, m2, t));
618 t2 = fold_binary (loop->cond_code, boolean_type_node,
619 n1cur, n2cur);
620 tree t3 = fold_binary (MULT_EXPR, itype, m1, ostep);
621 tree t4 = fold_binary (MULT_EXPR, itype, m2, ostep);
622 tree diff;
623 if (adj_first)
624 {
625 tree new_first;
626 if (integer_nonzerop (t2))
627 {
628 new_first = t;
629 n1first = n1cur;
630 n2first = n2cur;
631 if (flag_checking)
632 {
633 t3 = fold_binary (MINUS_EXPR, itype, n1cur, t3);
634 t4 = fold_binary (MINUS_EXPR, itype, n2cur, t4);
635 t3 = fold_binary (loop->cond_code,
636 boolean_type_node, t3, t4);
637 gcc_assert (integer_zerop (t3));
638 }
639 }
640 else
641 {
642 t3 = fold_binary (PLUS_EXPR, itype, n1cur, t3);
643 t4 = fold_binary (PLUS_EXPR, itype, n2cur, t4);
644 new_first = fold_binary (PLUS_EXPR, itype, t, ostep);
645 n1first = t3;
646 n2first = t4;
647 if (flag_checking)
648 {
649 t3 = fold_binary (loop->cond_code,
650 boolean_type_node, t3, t4);
651 gcc_assert (integer_nonzerop (t3));
652 }
653 }
654 diff = fold_binary (MINUS_EXPR, itype, new_first, first);
655 first = new_first;
656 fd->adjn1 = first;
657 }
658 else
659 {
660 tree new_last;
661 if (integer_zerop (t2))
662 {
663 t3 = fold_binary (MINUS_EXPR, itype, n1cur, t3);
664 t4 = fold_binary (MINUS_EXPR, itype, n2cur, t4);
665 new_last = fold_binary (MINUS_EXPR, itype, t, ostep);
666 n1last = t3;
667 n2last = t4;
668 if (flag_checking)
669 {
670 t3 = fold_binary (loop->cond_code,
671 boolean_type_node, t3, t4);
672 gcc_assert (integer_nonzerop (t3));
673 }
674 }
675 else
676 {
677 new_last = t;
678 n1last = n1cur;
679 n2last = n2cur;
680 if (flag_checking)
681 {
682 t3 = fold_binary (PLUS_EXPR, itype, n1cur, t3);
683 t4 = fold_binary (PLUS_EXPR, itype, n2cur, t4);
684 t3 = fold_binary (loop->cond_code,
685 boolean_type_node, t3, t4);
686 gcc_assert (integer_zerop (t3));
687 }
688 }
689 diff = fold_binary (MINUS_EXPR, itype, last, new_last);
690 }
691 if (TYPE_UNSIGNED (itype)
692 && single_nonrect_cond_code == GT_EXPR)
693 diff = fold_binary (TRUNC_DIV_EXPR, itype,
694 fold_unary (NEGATE_EXPR, itype, diff),
695 fold_unary (NEGATE_EXPR, itype,
696 ostep));
697 else
698 diff = fold_binary (TRUNC_DIV_EXPR, itype, diff, ostep);
699 diff = fold_convert (long_long_unsigned_type_node, diff);
700 single_nonrect_count
701 = fold_binary (MINUS_EXPR, long_long_unsigned_type_node,
702 single_nonrect_count, diff);
703 t = NULL_TREE;
704 }
705 }
706 else
707 t = fold_binary (loop->cond_code, boolean_type_node,
708 fold_convert (TREE_TYPE (loop->v), loop->n1),
709 fold_convert (TREE_TYPE (loop->v), loop->n2));
710 if (t && integer_zerop (t))
711 count = build_zero_cst (long_long_unsigned_type_node);
712 else if ((i == 0 || count != NULL_TREE)
713 && TREE_CODE (TREE_TYPE (loop->v)) == INTEGER_TYPE
714 && TREE_CONSTANT (loop->n1)
715 && TREE_CONSTANT (loop->n2)
716 && TREE_CODE (loop->step) == INTEGER_CST)
717 {
718 tree itype = TREE_TYPE (loop->v);
719
720 if (POINTER_TYPE_P (itype))
721 itype = signed_type_for (itype);
722 t = build_int_cst (itype, (loop->cond_code == LT_EXPR ? -1 : 1));
723 t = fold_build2 (PLUS_EXPR, itype,
724 fold_convert (itype, loop->step), t);
725 tree n1 = loop->n1;
726 tree n2 = loop->n2;
727 if (loop->m1 || loop->m2)
728 {
729 gcc_assert (single_nonrect != -1);
730 n1 = n1first;
731 n2 = n2first;
732 }
733 t = fold_build2 (PLUS_EXPR, itype, t, fold_convert (itype, n2));
734 t = fold_build2 (MINUS_EXPR, itype, t, fold_convert (itype, n1));
735 tree step = fold_convert_loc (loc, itype, loop->step);
736 if (TYPE_UNSIGNED (itype) && loop->cond_code == GT_EXPR)
737 t = fold_build2 (TRUNC_DIV_EXPR, itype,
738 fold_build1 (NEGATE_EXPR, itype, t),
739 fold_build1 (NEGATE_EXPR, itype, step));
740 else
741 t = fold_build2 (TRUNC_DIV_EXPR, itype, t, step);
742 tree llutype = long_long_unsigned_type_node;
743 t = fold_convert (llutype, t);
744 if (loop->m1 || loop->m2)
745 {
746 /* t is number of iterations of inner loop at either first
747 or last value of the outer iterator (the one with fewer
748 iterations).
749 Compute t2 = ((m2 - m1) * ostep) / step
750 and niters = outer_count * t
751 + t2 * ((outer_count - 1) * outer_count / 2)
752 */
753 tree m1 = loop->m1 ? loop->m1 : integer_zero_node;
754 tree m2 = loop->m2 ? loop->m2 : integer_zero_node;
755 m1 = fold_convert (itype, m1);
756 m2 = fold_convert (itype, m2);
757 tree t2 = fold_build2 (MINUS_EXPR, itype, m2, m1);
758 t2 = fold_build2 (MULT_EXPR, itype, t2, ostep);
759 if (TYPE_UNSIGNED (itype) && loop->cond_code == GT_EXPR)
760 t2 = fold_build2 (TRUNC_DIV_EXPR, itype,
761 fold_build1 (NEGATE_EXPR, itype, t2),
762 fold_build1 (NEGATE_EXPR, itype, step));
763 else
764 t2 = fold_build2 (TRUNC_DIV_EXPR, itype, t2, step);
765 t2 = fold_convert (llutype, t2);
766 fd->first_inner_iterations = t;
767 fd->factor = t2;
768 t = fold_build2 (MULT_EXPR, llutype, t,
769 single_nonrect_count);
770 tree t3 = fold_build2 (MINUS_EXPR, llutype,
771 single_nonrect_count,
772 build_one_cst (llutype));
773 t3 = fold_build2 (MULT_EXPR, llutype, t3,
774 single_nonrect_count);
775 t3 = fold_build2 (TRUNC_DIV_EXPR, llutype, t3,
776 build_int_cst (llutype, 2));
777 t2 = fold_build2 (MULT_EXPR, llutype, t2, t3);
778 t = fold_build2 (PLUS_EXPR, llutype, t, t2);
779 }
780 if (i == single_nonrect)
781 {
782 if (integer_zerop (t) || TREE_CODE (t) != INTEGER_CST)
783 count = t;
784 else
785 {
786 single_nonrect_count = t;
787 single_nonrect_cond_code = loop->cond_code;
788 if (count == NULL_TREE)
789 count = build_one_cst (llutype);
790 }
791 }
792 else if (count != NULL_TREE)
793 count = fold_build2 (MULT_EXPR, llutype, count, t);
794 else
795 count = t;
796 if (TREE_CODE (count) != INTEGER_CST)
797 count = NULL_TREE;
798 }
799 else if (count && !integer_zerop (count))
800 count = NULL_TREE;
801 }
802 }
803
804 if (count
805 && !simd
806 && (fd->sched_kind != OMP_CLAUSE_SCHEDULE_STATIC
807 || fd->have_ordered))
808 {
809 if (!tree_int_cst_lt (count, TYPE_MAX_VALUE (long_integer_type_node)))
810 iter_type = long_long_unsigned_type_node;
811 else
812 iter_type = long_integer_type_node;
813 }
814 else if (collapse_iter && *collapse_iter != NULL)
815 iter_type = TREE_TYPE (*collapse_iter);
816 fd->iter_type = iter_type;
817 if (collapse_iter && *collapse_iter == NULL)
818 *collapse_iter = create_tmp_var (iter_type, ".iter");
819 if (collapse_count && *collapse_count == NULL)
820 {
821 if (count)
822 {
823 *collapse_count = fold_convert_loc (loc, iter_type, count);
824 if (fd->first_inner_iterations && fd->factor)
825 {
826 t = make_tree_vec (4);
827 TREE_VEC_ELT (t, 0) = *collapse_count;
828 TREE_VEC_ELT (t, 1) = fd->first_inner_iterations;
829 TREE_VEC_ELT (t, 2) = fd->factor;
830 TREE_VEC_ELT (t, 3) = fd->adjn1;
831 *collapse_count = t;
832 }
833 }
834 else
835 *collapse_count = create_tmp_var (iter_type, ".count");
836 }
837
838 if (fd->collapse > 1 || fd->tiling || (fd->ordered && loops))
839 {
840 fd->loop.v = *collapse_iter;
841 fd->loop.n1 = build_int_cst (TREE_TYPE (fd->loop.v), 0);
842 fd->loop.n2 = *collapse_count;
843 if (TREE_CODE (fd->loop.n2) == TREE_VEC)
844 {
845 gcc_assert (fd->non_rect);
846 fd->first_inner_iterations = TREE_VEC_ELT (fd->loop.n2, 1);
847 fd->factor = TREE_VEC_ELT (fd->loop.n2, 2);
848 fd->adjn1 = TREE_VEC_ELT (fd->loop.n2, 3);
849 fd->loop.n2 = TREE_VEC_ELT (fd->loop.n2, 0);
850 }
851 fd->loop.step = build_int_cst (TREE_TYPE (fd->loop.v), 1);
852 fd->loop.m1 = NULL_TREE;
853 fd->loop.m2 = NULL_TREE;
854 fd->loop.outer = 0;
855 fd->loop.cond_code = LT_EXPR;
856 }
857 else if (loops)
858 loops[0] = fd->loop;
859 }
860
861 /* Build a call to GOMP_barrier. */
862
863 gimple *
864 omp_build_barrier (tree lhs)
865 {
866 tree fndecl = builtin_decl_explicit (lhs ? BUILT_IN_GOMP_BARRIER_CANCEL
867 : BUILT_IN_GOMP_BARRIER);
868 gcall *g = gimple_build_call (fndecl, 0);
869 if (lhs)
870 gimple_call_set_lhs (g, lhs);
871 return g;
872 }
873
874 /* Find OMP_FOR resp. OMP_SIMD with non-NULL OMP_FOR_INIT. Also, fill in pdata
875 array, pdata[0] non-NULL if there is anything non-trivial in between,
876 pdata[1] is address of OMP_PARALLEL in between if any, pdata[2] is address
877 of OMP_FOR in between if any and pdata[3] is address of the inner
878 OMP_FOR/OMP_SIMD. */
879
880 tree
881 find_combined_omp_for (tree *tp, int *walk_subtrees, void *data)
882 {
883 tree **pdata = (tree **) data;
884 *walk_subtrees = 0;
885 switch (TREE_CODE (*tp))
886 {
887 case OMP_FOR:
888 if (OMP_FOR_INIT (*tp) != NULL_TREE)
889 {
890 pdata[3] = tp;
891 return *tp;
892 }
893 pdata[2] = tp;
894 *walk_subtrees = 1;
895 break;
896 case OMP_SIMD:
897 if (OMP_FOR_INIT (*tp) != NULL_TREE)
898 {
899 pdata[3] = tp;
900 return *tp;
901 }
902 break;
903 case BIND_EXPR:
904 if (BIND_EXPR_VARS (*tp)
905 || (BIND_EXPR_BLOCK (*tp)
906 && BLOCK_VARS (BIND_EXPR_BLOCK (*tp))))
907 pdata[0] = tp;
908 *walk_subtrees = 1;
909 break;
910 case STATEMENT_LIST:
911 if (!tsi_one_before_end_p (tsi_start (*tp)))
912 pdata[0] = tp;
913 *walk_subtrees = 1;
914 break;
915 case TRY_FINALLY_EXPR:
916 pdata[0] = tp;
917 *walk_subtrees = 1;
918 break;
919 case OMP_PARALLEL:
920 pdata[1] = tp;
921 *walk_subtrees = 1;
922 break;
923 default:
924 break;
925 }
926 return NULL_TREE;
927 }
928
929 /* Return maximum possible vectorization factor for the target. */
930
931 poly_uint64
932 omp_max_vf (void)
933 {
934 if (!optimize
935 || optimize_debug
936 || !flag_tree_loop_optimize
937 || (!flag_tree_loop_vectorize
938 && global_options_set.x_flag_tree_loop_vectorize))
939 return 1;
940
941 auto_vector_modes modes;
942 targetm.vectorize.autovectorize_vector_modes (&modes, true);
943 if (!modes.is_empty ())
944 {
945 poly_uint64 vf = 0;
946 for (unsigned int i = 0; i < modes.length (); ++i)
947 /* The returned modes use the smallest element size (and thus
948 the largest nunits) for the vectorization approach that they
949 represent. */
950 vf = ordered_max (vf, GET_MODE_NUNITS (modes[i]));
951 return vf;
952 }
953
954 machine_mode vqimode = targetm.vectorize.preferred_simd_mode (QImode);
955 if (GET_MODE_CLASS (vqimode) == MODE_VECTOR_INT)
956 return GET_MODE_NUNITS (vqimode);
957
958 return 1;
959 }
960
961 /* Return maximum SIMT width if offloading may target SIMT hardware. */
962
963 int
964 omp_max_simt_vf (void)
965 {
966 if (!optimize)
967 return 0;
968 if (ENABLE_OFFLOADING)
969 for (const char *c = getenv ("OFFLOAD_TARGET_NAMES"); c;)
970 {
971 if (!strncmp (c, "nvptx", strlen ("nvptx")))
972 return 32;
973 else if ((c = strchr (c, ':')))
974 c++;
975 }
976 return 0;
977 }
978
979 /* Store the construct selectors as tree codes from last to first,
980 return their number. */
981
982 int
983 omp_constructor_traits_to_codes (tree ctx, enum tree_code *constructs)
984 {
985 int nconstructs = list_length (ctx);
986 int i = nconstructs - 1;
987 for (tree t2 = ctx; t2; t2 = TREE_CHAIN (t2), i--)
988 {
989 const char *sel = IDENTIFIER_POINTER (TREE_PURPOSE (t2));
990 if (!strcmp (sel, "target"))
991 constructs[i] = OMP_TARGET;
992 else if (!strcmp (sel, "teams"))
993 constructs[i] = OMP_TEAMS;
994 else if (!strcmp (sel, "parallel"))
995 constructs[i] = OMP_PARALLEL;
996 else if (!strcmp (sel, "for") || !strcmp (sel, "do"))
997 constructs[i] = OMP_FOR;
998 else if (!strcmp (sel, "simd"))
999 constructs[i] = OMP_SIMD;
1000 else
1001 gcc_unreachable ();
1002 }
1003 gcc_assert (i == -1);
1004 return nconstructs;
1005 }
1006
1007 /* Return true if PROP is possibly present in one of the offloading target's
1008 OpenMP contexts. The format of PROPS string is always offloading target's
1009 name terminated by '\0', followed by properties for that offloading
1010 target separated by '\0' and terminated by another '\0'. The strings
1011 are created from omp-device-properties installed files of all configured
1012 offloading targets. */
1013
1014 static bool
1015 omp_offload_device_kind_arch_isa (const char *props, const char *prop)
1016 {
1017 const char *names = getenv ("OFFLOAD_TARGET_NAMES");
1018 if (names == NULL || *names == '\0')
1019 return false;
1020 while (*props != '\0')
1021 {
1022 size_t name_len = strlen (props);
1023 bool matches = false;
1024 for (const char *c = names; c; )
1025 {
1026 if (strncmp (props, c, name_len) == 0
1027 && (c[name_len] == '\0'
1028 || c[name_len] == ':'
1029 || c[name_len] == '='))
1030 {
1031 matches = true;
1032 break;
1033 }
1034 else if ((c = strchr (c, ':')))
1035 c++;
1036 }
1037 props = props + name_len + 1;
1038 while (*props != '\0')
1039 {
1040 if (matches && strcmp (props, prop) == 0)
1041 return true;
1042 props = strchr (props, '\0') + 1;
1043 }
1044 props++;
1045 }
1046 return false;
1047 }
1048
1049 /* Return true if the current code location is or might be offloaded.
1050 Return true in declare target functions, or when nested in a target
1051 region or when unsure, return false otherwise. */
1052
1053 static bool
1054 omp_maybe_offloaded (void)
1055 {
1056 if (!ENABLE_OFFLOADING)
1057 return false;
1058 const char *names = getenv ("OFFLOAD_TARGET_NAMES");
1059 if (names == NULL || *names == '\0')
1060 return false;
1061
1062 if (symtab->state == PARSING)
1063 /* Maybe. */
1064 return true;
1065 if (cfun && cfun->after_inlining)
1066 return false;
1067 if (current_function_decl
1068 && lookup_attribute ("omp declare target",
1069 DECL_ATTRIBUTES (current_function_decl)))
1070 return true;
1071 if (cfun && (cfun->curr_properties & PROP_gimple_any) == 0)
1072 {
1073 enum tree_code construct = OMP_TARGET;
1074 if (omp_construct_selector_matches (&construct, 1, NULL))
1075 return true;
1076 }
1077 return false;
1078 }
1079
1080 /* Return a name from PROP, a property in selectors accepting
1081 name lists. */
1082
1083 static const char *
1084 omp_context_name_list_prop (tree prop)
1085 {
1086 if (TREE_PURPOSE (prop))
1087 return IDENTIFIER_POINTER (TREE_PURPOSE (prop));
1088 else
1089 {
1090 const char *ret = TREE_STRING_POINTER (TREE_VALUE (prop));
1091 if ((size_t) TREE_STRING_LENGTH (TREE_VALUE (prop)) == strlen (ret) + 1)
1092 return ret;
1093 return NULL;
1094 }
1095 }
1096
1097 /* Return 1 if context selector matches the current OpenMP context, 0
1098 if it does not and -1 if it is unknown and need to be determined later.
1099 Some properties can be checked right away during parsing (this routine),
1100 others need to wait until the whole TU is parsed, others need to wait until
1101 IPA, others until vectorization. */
1102
1103 int
1104 omp_context_selector_matches (tree ctx)
1105 {
1106 int ret = 1;
1107 for (tree t1 = ctx; t1; t1 = TREE_CHAIN (t1))
1108 {
1109 char set = IDENTIFIER_POINTER (TREE_PURPOSE (t1))[0];
1110 if (set == 'c')
1111 {
1112 /* For now, ignore the construct set. While something can be
1113 determined already during parsing, we don't know until end of TU
1114 whether additional constructs aren't added through declare variant
1115 unless "omp declare variant variant" attribute exists already
1116 (so in most of the cases), and we'd need to maintain set of
1117 surrounding OpenMP constructs, which is better handled during
1118 gimplification. */
1119 if (symtab->state == PARSING)
1120 {
1121 ret = -1;
1122 continue;
1123 }
1124
1125 enum tree_code constructs[5];
1126 int nconstructs
1127 = omp_constructor_traits_to_codes (TREE_VALUE (t1), constructs);
1128
1129 if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
1130 {
1131 if (!cfun->after_inlining)
1132 {
1133 ret = -1;
1134 continue;
1135 }
1136 int i;
1137 for (i = 0; i < nconstructs; ++i)
1138 if (constructs[i] == OMP_SIMD)
1139 break;
1140 if (i < nconstructs)
1141 {
1142 ret = -1;
1143 continue;
1144 }
1145 /* If there is no simd, assume it is ok after IPA,
1146 constructs should have been checked before. */
1147 continue;
1148 }
1149
1150 int r = omp_construct_selector_matches (constructs, nconstructs,
1151 NULL);
1152 if (r == 0)
1153 return 0;
1154 if (r == -1)
1155 ret = -1;
1156 continue;
1157 }
1158 for (tree t2 = TREE_VALUE (t1); t2; t2 = TREE_CHAIN (t2))
1159 {
1160 const char *sel = IDENTIFIER_POINTER (TREE_PURPOSE (t2));
1161 switch (*sel)
1162 {
1163 case 'v':
1164 if (set == 'i' && !strcmp (sel, "vendor"))
1165 for (tree t3 = TREE_VALUE (t2); t3; t3 = TREE_CHAIN (t3))
1166 {
1167 const char *prop = omp_context_name_list_prop (t3);
1168 if (prop == NULL)
1169 return 0;
1170 if ((!strcmp (prop, " score") && TREE_PURPOSE (t3))
1171 || !strcmp (prop, "gnu"))
1172 continue;
1173 return 0;
1174 }
1175 break;
1176 case 'e':
1177 if (set == 'i' && !strcmp (sel, "extension"))
1178 /* We don't support any extensions right now. */
1179 return 0;
1180 break;
1181 case 'a':
1182 if (set == 'i' && !strcmp (sel, "atomic_default_mem_order"))
1183 {
1184 if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
1185 break;
1186
1187 enum omp_memory_order omo
1188 = ((enum omp_memory_order)
1189 (omp_requires_mask
1190 & OMP_REQUIRES_ATOMIC_DEFAULT_MEM_ORDER));
1191 if (omo == OMP_MEMORY_ORDER_UNSPECIFIED)
1192 {
1193 /* We don't know yet, until end of TU. */
1194 if (symtab->state == PARSING)
1195 {
1196 ret = -1;
1197 break;
1198 }
1199 else
1200 omo = OMP_MEMORY_ORDER_RELAXED;
1201 }
1202 tree t3 = TREE_VALUE (t2);
1203 const char *prop = IDENTIFIER_POINTER (TREE_PURPOSE (t3));
1204 if (!strcmp (prop, " score"))
1205 {
1206 t3 = TREE_CHAIN (t3);
1207 prop = IDENTIFIER_POINTER (TREE_PURPOSE (t3));
1208 }
1209 if (!strcmp (prop, "relaxed")
1210 && omo != OMP_MEMORY_ORDER_RELAXED)
1211 return 0;
1212 else if (!strcmp (prop, "seq_cst")
1213 && omo != OMP_MEMORY_ORDER_SEQ_CST)
1214 return 0;
1215 else if (!strcmp (prop, "acq_rel")
1216 && omo != OMP_MEMORY_ORDER_ACQ_REL)
1217 return 0;
1218 }
1219 if (set == 'd' && !strcmp (sel, "arch"))
1220 for (tree t3 = TREE_VALUE (t2); t3; t3 = TREE_CHAIN (t3))
1221 {
1222 const char *arch = omp_context_name_list_prop (t3);
1223 if (arch == NULL)
1224 return 0;
1225 int r = 0;
1226 if (targetm.omp.device_kind_arch_isa != NULL)
1227 r = targetm.omp.device_kind_arch_isa (omp_device_arch,
1228 arch);
1229 if (r == 0 || (r == -1 && symtab->state != PARSING))
1230 {
1231 /* If we are or might be in a target region or
1232 declare target function, need to take into account
1233 also offloading values. */
1234 if (!omp_maybe_offloaded ())
1235 return 0;
1236 if (ENABLE_OFFLOADING)
1237 {
1238 const char *arches = omp_offload_device_arch;
1239 if (omp_offload_device_kind_arch_isa (arches,
1240 arch))
1241 {
1242 ret = -1;
1243 continue;
1244 }
1245 }
1246 return 0;
1247 }
1248 else if (r == -1)
1249 ret = -1;
1250 /* If arch matches on the host, it still might not match
1251 in the offloading region. */
1252 else if (omp_maybe_offloaded ())
1253 ret = -1;
1254 }
1255 break;
1256 case 'u':
1257 if (set == 'i' && !strcmp (sel, "unified_address"))
1258 {
1259 if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
1260 break;
1261
1262 if ((omp_requires_mask & OMP_REQUIRES_UNIFIED_ADDRESS) == 0)
1263 {
1264 if (symtab->state == PARSING)
1265 ret = -1;
1266 else
1267 return 0;
1268 }
1269 break;
1270 }
1271 if (set == 'i' && !strcmp (sel, "unified_shared_memory"))
1272 {
1273 if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
1274 break;
1275
1276 if ((omp_requires_mask
1277 & OMP_REQUIRES_UNIFIED_SHARED_MEMORY) == 0)
1278 {
1279 if (symtab->state == PARSING)
1280 ret = -1;
1281 else
1282 return 0;
1283 }
1284 break;
1285 }
1286 break;
1287 case 'd':
1288 if (set == 'i' && !strcmp (sel, "dynamic_allocators"))
1289 {
1290 if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
1291 break;
1292
1293 if ((omp_requires_mask
1294 & OMP_REQUIRES_DYNAMIC_ALLOCATORS) == 0)
1295 {
1296 if (symtab->state == PARSING)
1297 ret = -1;
1298 else
1299 return 0;
1300 }
1301 break;
1302 }
1303 break;
1304 case 'r':
1305 if (set == 'i' && !strcmp (sel, "reverse_offload"))
1306 {
1307 if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
1308 break;
1309
1310 if ((omp_requires_mask & OMP_REQUIRES_REVERSE_OFFLOAD) == 0)
1311 {
1312 if (symtab->state == PARSING)
1313 ret = -1;
1314 else
1315 return 0;
1316 }
1317 break;
1318 }
1319 break;
1320 case 'k':
1321 if (set == 'd' && !strcmp (sel, "kind"))
1322 for (tree t3 = TREE_VALUE (t2); t3; t3 = TREE_CHAIN (t3))
1323 {
1324 const char *prop = omp_context_name_list_prop (t3);
1325 if (prop == NULL)
1326 return 0;
1327 if (!strcmp (prop, "any"))
1328 continue;
1329 if (!strcmp (prop, "host"))
1330 {
1331 if (omp_maybe_offloaded ())
1332 ret = -1;
1333 continue;
1334 }
1335 if (!strcmp (prop, "nohost"))
1336 {
1337 if (omp_maybe_offloaded ())
1338 ret = -1;
1339 else
1340 return 0;
1341 continue;
1342 }
1343 int r = 0;
1344 if (targetm.omp.device_kind_arch_isa != NULL)
1345 r = targetm.omp.device_kind_arch_isa (omp_device_kind,
1346 prop);
1347 else
1348 r = strcmp (prop, "cpu") == 0;
1349 if (r == 0 || (r == -1 && symtab->state != PARSING))
1350 {
1351 /* If we are or might be in a target region or
1352 declare target function, need to take into account
1353 also offloading values. */
1354 if (!omp_maybe_offloaded ())
1355 return 0;
1356 if (ENABLE_OFFLOADING)
1357 {
1358 const char *kinds = omp_offload_device_kind;
1359 if (omp_offload_device_kind_arch_isa (kinds, prop))
1360 {
1361 ret = -1;
1362 continue;
1363 }
1364 }
1365 return 0;
1366 }
1367 else if (r == -1)
1368 ret = -1;
1369 /* If kind matches on the host, it still might not match
1370 in the offloading region. */
1371 else if (omp_maybe_offloaded ())
1372 ret = -1;
1373 }
1374 break;
1375 case 'i':
1376 if (set == 'd' && !strcmp (sel, "isa"))
1377 for (tree t3 = TREE_VALUE (t2); t3; t3 = TREE_CHAIN (t3))
1378 {
1379 const char *isa = omp_context_name_list_prop (t3);
1380 if (isa == NULL)
1381 return 0;
1382 int r = 0;
1383 if (targetm.omp.device_kind_arch_isa != NULL)
1384 r = targetm.omp.device_kind_arch_isa (omp_device_isa,
1385 isa);
1386 if (r == 0 || (r == -1 && symtab->state != PARSING))
1387 {
1388 /* If isa is valid on the target, but not in the
1389 current function and current function has
1390 #pragma omp declare simd on it, some simd clones
1391 might have the isa added later on. */
1392 if (r == -1
1393 && targetm.simd_clone.compute_vecsize_and_simdlen
1394 && (cfun == NULL || !cfun->after_inlining))
1395 {
1396 tree attrs
1397 = DECL_ATTRIBUTES (current_function_decl);
1398 if (lookup_attribute ("omp declare simd", attrs))
1399 {
1400 ret = -1;
1401 continue;
1402 }
1403 }
1404 /* If we are or might be in a target region or
1405 declare target function, need to take into account
1406 also offloading values. */
1407 if (!omp_maybe_offloaded ())
1408 return 0;
1409 if (ENABLE_OFFLOADING)
1410 {
1411 const char *isas = omp_offload_device_isa;
1412 if (omp_offload_device_kind_arch_isa (isas, isa))
1413 {
1414 ret = -1;
1415 continue;
1416 }
1417 }
1418 return 0;
1419 }
1420 else if (r == -1)
1421 ret = -1;
1422 /* If isa matches on the host, it still might not match
1423 in the offloading region. */
1424 else if (omp_maybe_offloaded ())
1425 ret = -1;
1426 }
1427 break;
1428 case 'c':
1429 if (set == 'u' && !strcmp (sel, "condition"))
1430 for (tree t3 = TREE_VALUE (t2); t3; t3 = TREE_CHAIN (t3))
1431 if (TREE_PURPOSE (t3) == NULL_TREE)
1432 {
1433 if (integer_zerop (TREE_VALUE (t3)))
1434 return 0;
1435 if (integer_nonzerop (TREE_VALUE (t3)))
1436 break;
1437 ret = -1;
1438 }
1439 break;
1440 default:
1441 break;
1442 }
1443 }
1444 }
1445 return ret;
1446 }
1447
1448 /* Compare construct={simd} CLAUSES1 with CLAUSES2, return 0/-1/1/2 as
1449 in omp_context_selector_set_compare. */
1450
1451 static int
1452 omp_construct_simd_compare (tree clauses1, tree clauses2)
1453 {
1454 if (clauses1 == NULL_TREE)
1455 return clauses2 == NULL_TREE ? 0 : -1;
1456 if (clauses2 == NULL_TREE)
1457 return 1;
1458
1459 int r = 0;
1460 struct declare_variant_simd_data {
1461 bool inbranch, notinbranch;
1462 tree simdlen;
1463 auto_vec<tree,16> data_sharing;
1464 auto_vec<tree,16> aligned;
1465 declare_variant_simd_data ()
1466 : inbranch(false), notinbranch(false), simdlen(NULL_TREE) {}
1467 } data[2];
1468 unsigned int i;
1469 for (i = 0; i < 2; i++)
1470 for (tree c = i ? clauses2 : clauses1; c; c = OMP_CLAUSE_CHAIN (c))
1471 {
1472 vec<tree> *v;
1473 switch (OMP_CLAUSE_CODE (c))
1474 {
1475 case OMP_CLAUSE_INBRANCH:
1476 data[i].inbranch = true;
1477 continue;
1478 case OMP_CLAUSE_NOTINBRANCH:
1479 data[i].notinbranch = true;
1480 continue;
1481 case OMP_CLAUSE_SIMDLEN:
1482 data[i].simdlen = OMP_CLAUSE_SIMDLEN_EXPR (c);
1483 continue;
1484 case OMP_CLAUSE_UNIFORM:
1485 case OMP_CLAUSE_LINEAR:
1486 v = &data[i].data_sharing;
1487 break;
1488 case OMP_CLAUSE_ALIGNED:
1489 v = &data[i].aligned;
1490 break;
1491 default:
1492 gcc_unreachable ();
1493 }
1494 unsigned HOST_WIDE_INT argno = tree_to_uhwi (OMP_CLAUSE_DECL (c));
1495 if (argno >= v->length ())
1496 v->safe_grow_cleared (argno + 1, true);
1497 (*v)[argno] = c;
1498 }
1499 /* Here, r is used as a bitmask, 2 is set if CLAUSES1 has something
1500 CLAUSES2 doesn't, 1 is set if CLAUSES2 has something CLAUSES1
1501 doesn't. Thus, r == 3 implies return value 2, r == 1 implies
1502 -1, r == 2 implies 1 and r == 0 implies 0. */
1503 if (data[0].inbranch != data[1].inbranch)
1504 r |= data[0].inbranch ? 2 : 1;
1505 if (data[0].notinbranch != data[1].notinbranch)
1506 r |= data[0].notinbranch ? 2 : 1;
1507 if (!simple_cst_equal (data[0].simdlen, data[1].simdlen))
1508 {
1509 if (data[0].simdlen && data[1].simdlen)
1510 return 2;
1511 r |= data[0].simdlen ? 2 : 1;
1512 }
1513 if (data[0].data_sharing.length () < data[1].data_sharing.length ()
1514 || data[0].aligned.length () < data[1].aligned.length ())
1515 r |= 1;
1516 tree c1, c2;
1517 FOR_EACH_VEC_ELT (data[0].data_sharing, i, c1)
1518 {
1519 c2 = (i < data[1].data_sharing.length ()
1520 ? data[1].data_sharing[i] : NULL_TREE);
1521 if ((c1 == NULL_TREE) != (c2 == NULL_TREE))
1522 {
1523 r |= c1 != NULL_TREE ? 2 : 1;
1524 continue;
1525 }
1526 if (c1 == NULL_TREE)
1527 continue;
1528 if (OMP_CLAUSE_CODE (c1) != OMP_CLAUSE_CODE (c2))
1529 return 2;
1530 if (OMP_CLAUSE_CODE (c1) != OMP_CLAUSE_LINEAR)
1531 continue;
1532 if (OMP_CLAUSE_LINEAR_VARIABLE_STRIDE (c1)
1533 != OMP_CLAUSE_LINEAR_VARIABLE_STRIDE (c2))
1534 return 2;
1535 if (OMP_CLAUSE_LINEAR_KIND (c1) != OMP_CLAUSE_LINEAR_KIND (c2))
1536 return 2;
1537 if (!simple_cst_equal (OMP_CLAUSE_LINEAR_STEP (c1),
1538 OMP_CLAUSE_LINEAR_STEP (c2)))
1539 return 2;
1540 }
1541 FOR_EACH_VEC_ELT (data[0].aligned, i, c1)
1542 {
1543 c2 = i < data[1].aligned.length () ? data[1].aligned[i] : NULL_TREE;
1544 if ((c1 == NULL_TREE) != (c2 == NULL_TREE))
1545 {
1546 r |= c1 != NULL_TREE ? 2 : 1;
1547 continue;
1548 }
1549 if (c1 == NULL_TREE)
1550 continue;
1551 if (!simple_cst_equal (OMP_CLAUSE_ALIGNED_ALIGNMENT (c1),
1552 OMP_CLAUSE_ALIGNED_ALIGNMENT (c2)))
1553 return 2;
1554 }
1555 switch (r)
1556 {
1557 case 0: return 0;
1558 case 1: return -1;
1559 case 2: return 1;
1560 case 3: return 2;
1561 default: gcc_unreachable ();
1562 }
1563 }
1564
1565 /* Compare properties of selectors SEL from SET other than construct.
1566 Return 0/-1/1/2 as in omp_context_selector_set_compare.
1567 Unlike set names or selector names, properties can have duplicates. */
1568
1569 static int
1570 omp_context_selector_props_compare (const char *set, const char *sel,
1571 tree ctx1, tree ctx2)
1572 {
1573 int ret = 0;
1574 for (int pass = 0; pass < 2; pass++)
1575 for (tree t1 = pass ? ctx2 : ctx1; t1; t1 = TREE_CHAIN (t1))
1576 {
1577 tree t2;
1578 for (t2 = pass ? ctx1 : ctx2; t2; t2 = TREE_CHAIN (t2))
1579 if (TREE_PURPOSE (t1) == TREE_PURPOSE (t2))
1580 {
1581 if (TREE_PURPOSE (t1) == NULL_TREE)
1582 {
1583 if (set[0] == 'u' && strcmp (sel, "condition") == 0)
1584 {
1585 if (integer_zerop (TREE_VALUE (t1))
1586 != integer_zerop (TREE_VALUE (t2)))
1587 return 2;
1588 break;
1589 }
1590 if (simple_cst_equal (TREE_VALUE (t1), TREE_VALUE (t2)))
1591 break;
1592 }
1593 else if (strcmp (IDENTIFIER_POINTER (TREE_PURPOSE (t1)),
1594 " score") == 0)
1595 {
1596 if (!simple_cst_equal (TREE_VALUE (t1), TREE_VALUE (t2)))
1597 return 2;
1598 break;
1599 }
1600 else
1601 break;
1602 }
1603 else if (TREE_PURPOSE (t1)
1604 && TREE_PURPOSE (t2) == NULL_TREE
1605 && TREE_CODE (TREE_VALUE (t2)) == STRING_CST)
1606 {
1607 const char *p1 = omp_context_name_list_prop (t1);
1608 const char *p2 = omp_context_name_list_prop (t2);
1609 if (p2
1610 && strcmp (p1, p2) == 0
1611 && strcmp (p1, " score"))
1612 break;
1613 }
1614 else if (TREE_PURPOSE (t1) == NULL_TREE
1615 && TREE_PURPOSE (t2)
1616 && TREE_CODE (TREE_VALUE (t1)) == STRING_CST)
1617 {
1618 const char *p1 = omp_context_name_list_prop (t1);
1619 const char *p2 = omp_context_name_list_prop (t2);
1620 if (p1
1621 && strcmp (p1, p2) == 0
1622 && strcmp (p1, " score"))
1623 break;
1624 }
1625 if (t2 == NULL_TREE)
1626 {
1627 int r = pass ? -1 : 1;
1628 if (ret && ret != r)
1629 return 2;
1630 else if (pass)
1631 return r;
1632 else
1633 {
1634 ret = r;
1635 break;
1636 }
1637 }
1638 }
1639 return ret;
1640 }
1641
1642 /* Compare single context selector sets CTX1 and CTX2 with SET name.
1643 Return 0 if CTX1 is equal to CTX2,
1644 -1 if CTX1 is a strict subset of CTX2,
1645 1 if CTX2 is a strict subset of CTX1, or
1646 2 if neither context is a subset of another one. */
1647
1648 int
1649 omp_context_selector_set_compare (const char *set, tree ctx1, tree ctx2)
1650 {
1651 bool swapped = false;
1652 int ret = 0;
1653 int len1 = list_length (ctx1);
1654 int len2 = list_length (ctx2);
1655 int cnt = 0;
1656 if (len1 < len2)
1657 {
1658 swapped = true;
1659 std::swap (ctx1, ctx2);
1660 std::swap (len1, len2);
1661 }
1662 if (set[0] == 'c')
1663 {
1664 tree t1;
1665 tree t2 = ctx2;
1666 tree simd = get_identifier ("simd");
1667 /* Handle construct set specially. In this case the order
1668 of the selector matters too. */
1669 for (t1 = ctx1; t1; t1 = TREE_CHAIN (t1))
1670 if (TREE_PURPOSE (t1) == TREE_PURPOSE (t2))
1671 {
1672 int r = 0;
1673 if (TREE_PURPOSE (t1) == simd)
1674 r = omp_construct_simd_compare (TREE_VALUE (t1),
1675 TREE_VALUE (t2));
1676 if (r == 2 || (ret && r && (ret < 0) != (r < 0)))
1677 return 2;
1678 if (ret == 0)
1679 ret = r;
1680 t2 = TREE_CHAIN (t2);
1681 if (t2 == NULL_TREE)
1682 {
1683 t1 = TREE_CHAIN (t1);
1684 break;
1685 }
1686 }
1687 else if (ret < 0)
1688 return 2;
1689 else
1690 ret = 1;
1691 if (t2 != NULL_TREE)
1692 return 2;
1693 if (t1 != NULL_TREE)
1694 {
1695 if (ret < 0)
1696 return 2;
1697 ret = 1;
1698 }
1699 if (ret == 0)
1700 return 0;
1701 return swapped ? -ret : ret;
1702 }
1703 for (tree t1 = ctx1; t1; t1 = TREE_CHAIN (t1))
1704 {
1705 tree t2;
1706 for (t2 = ctx2; t2; t2 = TREE_CHAIN (t2))
1707 if (TREE_PURPOSE (t1) == TREE_PURPOSE (t2))
1708 {
1709 const char *sel = IDENTIFIER_POINTER (TREE_PURPOSE (t1));
1710 int r = omp_context_selector_props_compare (set, sel,
1711 TREE_VALUE (t1),
1712 TREE_VALUE (t2));
1713 if (r == 2 || (ret && r && (ret < 0) != (r < 0)))
1714 return 2;
1715 if (ret == 0)
1716 ret = r;
1717 cnt++;
1718 break;
1719 }
1720 if (t2 == NULL_TREE)
1721 {
1722 if (ret == -1)
1723 return 2;
1724 ret = 1;
1725 }
1726 }
1727 if (cnt < len2)
1728 return 2;
1729 if (ret == 0)
1730 return 0;
1731 return swapped ? -ret : ret;
1732 }
1733
1734 /* Compare whole context selector specification CTX1 and CTX2.
1735 Return 0 if CTX1 is equal to CTX2,
1736 -1 if CTX1 is a strict subset of CTX2,
1737 1 if CTX2 is a strict subset of CTX1, or
1738 2 if neither context is a subset of another one. */
1739
1740 static int
1741 omp_context_selector_compare (tree ctx1, tree ctx2)
1742 {
1743 bool swapped = false;
1744 int ret = 0;
1745 int len1 = list_length (ctx1);
1746 int len2 = list_length (ctx2);
1747 int cnt = 0;
1748 if (len1 < len2)
1749 {
1750 swapped = true;
1751 std::swap (ctx1, ctx2);
1752 std::swap (len1, len2);
1753 }
1754 for (tree t1 = ctx1; t1; t1 = TREE_CHAIN (t1))
1755 {
1756 tree t2;
1757 for (t2 = ctx2; t2; t2 = TREE_CHAIN (t2))
1758 if (TREE_PURPOSE (t1) == TREE_PURPOSE (t2))
1759 {
1760 const char *set = IDENTIFIER_POINTER (TREE_PURPOSE (t1));
1761 int r = omp_context_selector_set_compare (set, TREE_VALUE (t1),
1762 TREE_VALUE (t2));
1763 if (r == 2 || (ret && r && (ret < 0) != (r < 0)))
1764 return 2;
1765 if (ret == 0)
1766 ret = r;
1767 cnt++;
1768 break;
1769 }
1770 if (t2 == NULL_TREE)
1771 {
1772 if (ret == -1)
1773 return 2;
1774 ret = 1;
1775 }
1776 }
1777 if (cnt < len2)
1778 return 2;
1779 if (ret == 0)
1780 return 0;
1781 return swapped ? -ret : ret;
1782 }
1783
1784 /* From context selector CTX, return trait-selector with name SEL in
1785 trait-selector-set with name SET if any, or NULL_TREE if not found.
1786 If SEL is NULL, return the list of trait-selectors in SET. */
1787
1788 tree
1789 omp_get_context_selector (tree ctx, const char *set, const char *sel)
1790 {
1791 tree setid = get_identifier (set);
1792 tree selid = sel ? get_identifier (sel) : NULL_TREE;
1793 for (tree t1 = ctx; t1; t1 = TREE_CHAIN (t1))
1794 if (TREE_PURPOSE (t1) == setid)
1795 {
1796 if (sel == NULL)
1797 return TREE_VALUE (t1);
1798 for (tree t2 = TREE_VALUE (t1); t2; t2 = TREE_CHAIN (t2))
1799 if (TREE_PURPOSE (t2) == selid)
1800 return t2;
1801 }
1802 return NULL_TREE;
1803 }
1804
1805 /* Compute *SCORE for context selector CTX. Return true if the score
1806 would be different depending on whether it is a declare simd clone or
1807 not. DECLARE_SIMD should be true for the case when it would be
1808 a declare simd clone. */
1809
1810 static bool
1811 omp_context_compute_score (tree ctx, widest_int *score, bool declare_simd)
1812 {
1813 tree construct = omp_get_context_selector (ctx, "construct", NULL);
1814 bool has_kind = omp_get_context_selector (ctx, "device", "kind");
1815 bool has_arch = omp_get_context_selector (ctx, "device", "arch");
1816 bool has_isa = omp_get_context_selector (ctx, "device", "isa");
1817 bool ret = false;
1818 *score = 1;
1819 for (tree t1 = ctx; t1; t1 = TREE_CHAIN (t1))
1820 if (TREE_VALUE (t1) != construct)
1821 for (tree t2 = TREE_VALUE (t1); t2; t2 = TREE_CHAIN (t2))
1822 if (tree t3 = TREE_VALUE (t2))
1823 if (TREE_PURPOSE (t3)
1824 && strcmp (IDENTIFIER_POINTER (TREE_PURPOSE (t3)), " score") == 0
1825 && TREE_CODE (TREE_VALUE (t3)) == INTEGER_CST)
1826 *score += wi::to_widest (TREE_VALUE (t3));
1827 if (construct || has_kind || has_arch || has_isa)
1828 {
1829 int scores[12];
1830 enum tree_code constructs[5];
1831 int nconstructs = 0;
1832 if (construct)
1833 nconstructs = omp_constructor_traits_to_codes (construct, constructs);
1834 if (omp_construct_selector_matches (constructs, nconstructs, scores)
1835 == 2)
1836 ret = true;
1837 int b = declare_simd ? nconstructs + 1 : 0;
1838 if (scores[b + nconstructs] + 4U < score->get_precision ())
1839 {
1840 for (int n = 0; n < nconstructs; ++n)
1841 {
1842 if (scores[b + n] < 0)
1843 {
1844 *score = -1;
1845 return ret;
1846 }
1847 *score += wi::shifted_mask <widest_int> (scores[b + n], 1, false);
1848 }
1849 if (has_kind)
1850 *score += wi::shifted_mask <widest_int> (scores[b + nconstructs],
1851 1, false);
1852 if (has_arch)
1853 *score += wi::shifted_mask <widest_int> (scores[b + nconstructs] + 1,
1854 1, false);
1855 if (has_isa)
1856 *score += wi::shifted_mask <widest_int> (scores[b + nconstructs] + 2,
1857 1, false);
1858 }
1859 else /* FIXME: Implement this. */
1860 gcc_unreachable ();
1861 }
1862 return ret;
1863 }
1864
1865 /* Class describing a single variant. */
1866 struct GTY(()) omp_declare_variant_entry {
1867 /* NODE of the variant. */
1868 cgraph_node *variant;
1869 /* Score if not in declare simd clone. */
1870 widest_int score;
1871 /* Score if in declare simd clone. */
1872 widest_int score_in_declare_simd_clone;
1873 /* Context selector for the variant. */
1874 tree ctx;
1875 /* True if the context selector is known to match already. */
1876 bool matches;
1877 };
1878
1879 /* Class describing a function with variants. */
1880 struct GTY((for_user)) omp_declare_variant_base_entry {
1881 /* NODE of the base function. */
1882 cgraph_node *base;
1883 /* NODE of the artificial function created for the deferred variant
1884 resolution. */
1885 cgraph_node *node;
1886 /* Vector of the variants. */
1887 vec<omp_declare_variant_entry, va_gc> *variants;
1888 };
1889
1890 struct omp_declare_variant_hasher
1891 : ggc_ptr_hash<omp_declare_variant_base_entry> {
1892 static hashval_t hash (omp_declare_variant_base_entry *);
1893 static bool equal (omp_declare_variant_base_entry *,
1894 omp_declare_variant_base_entry *);
1895 };
1896
1897 hashval_t
1898 omp_declare_variant_hasher::hash (omp_declare_variant_base_entry *x)
1899 {
1900 inchash::hash hstate;
1901 hstate.add_int (DECL_UID (x->base->decl));
1902 hstate.add_int (x->variants->length ());
1903 omp_declare_variant_entry *variant;
1904 unsigned int i;
1905 FOR_EACH_VEC_SAFE_ELT (x->variants, i, variant)
1906 {
1907 hstate.add_int (DECL_UID (variant->variant->decl));
1908 hstate.add_wide_int (variant->score);
1909 hstate.add_wide_int (variant->score_in_declare_simd_clone);
1910 hstate.add_ptr (variant->ctx);
1911 hstate.add_int (variant->matches);
1912 }
1913 return hstate.end ();
1914 }
1915
1916 bool
1917 omp_declare_variant_hasher::equal (omp_declare_variant_base_entry *x,
1918 omp_declare_variant_base_entry *y)
1919 {
1920 if (x->base != y->base
1921 || x->variants->length () != y->variants->length ())
1922 return false;
1923 omp_declare_variant_entry *variant;
1924 unsigned int i;
1925 FOR_EACH_VEC_SAFE_ELT (x->variants, i, variant)
1926 if (variant->variant != (*y->variants)[i].variant
1927 || variant->score != (*y->variants)[i].score
1928 || (variant->score_in_declare_simd_clone
1929 != (*y->variants)[i].score_in_declare_simd_clone)
1930 || variant->ctx != (*y->variants)[i].ctx
1931 || variant->matches != (*y->variants)[i].matches)
1932 return false;
1933 return true;
1934 }
1935
1936 static GTY(()) hash_table<omp_declare_variant_hasher> *omp_declare_variants;
1937
1938 struct omp_declare_variant_alt_hasher
1939 : ggc_ptr_hash<omp_declare_variant_base_entry> {
1940 static hashval_t hash (omp_declare_variant_base_entry *);
1941 static bool equal (omp_declare_variant_base_entry *,
1942 omp_declare_variant_base_entry *);
1943 };
1944
1945 hashval_t
1946 omp_declare_variant_alt_hasher::hash (omp_declare_variant_base_entry *x)
1947 {
1948 return DECL_UID (x->node->decl);
1949 }
1950
1951 bool
1952 omp_declare_variant_alt_hasher::equal (omp_declare_variant_base_entry *x,
1953 omp_declare_variant_base_entry *y)
1954 {
1955 return x->node == y->node;
1956 }
1957
1958 static GTY(()) hash_table<omp_declare_variant_alt_hasher>
1959 *omp_declare_variant_alt;
1960
1961 /* Try to resolve declare variant after gimplification. */
1962
1963 static tree
1964 omp_resolve_late_declare_variant (tree alt)
1965 {
1966 cgraph_node *node = cgraph_node::get (alt);
1967 cgraph_node *cur_node = cgraph_node::get (cfun->decl);
1968 if (node == NULL
1969 || !node->declare_variant_alt
1970 || !cfun->after_inlining)
1971 return alt;
1972
1973 omp_declare_variant_base_entry entry;
1974 entry.base = NULL;
1975 entry.node = node;
1976 entry.variants = NULL;
1977 omp_declare_variant_base_entry *entryp
1978 = omp_declare_variant_alt->find_with_hash (&entry, DECL_UID (alt));
1979
1980 unsigned int i, j;
1981 omp_declare_variant_entry *varentry1, *varentry2;
1982 auto_vec <bool, 16> matches;
1983 unsigned int nmatches = 0;
1984 FOR_EACH_VEC_SAFE_ELT (entryp->variants, i, varentry1)
1985 {
1986 if (varentry1->matches)
1987 {
1988 /* This has been checked to be ok already. */
1989 matches.safe_push (true);
1990 nmatches++;
1991 continue;
1992 }
1993 switch (omp_context_selector_matches (varentry1->ctx))
1994 {
1995 case 0:
1996 matches.safe_push (false);
1997 break;
1998 case -1:
1999 return alt;
2000 default:
2001 matches.safe_push (true);
2002 nmatches++;
2003 break;
2004 }
2005 }
2006
2007 if (nmatches == 0)
2008 return entryp->base->decl;
2009
2010 /* A context selector that is a strict subset of another context selector
2011 has a score of zero. */
2012 FOR_EACH_VEC_SAFE_ELT (entryp->variants, i, varentry1)
2013 if (matches[i])
2014 {
2015 for (j = i + 1;
2016 vec_safe_iterate (entryp->variants, j, &varentry2); ++j)
2017 if (matches[j])
2018 {
2019 int r = omp_context_selector_compare (varentry1->ctx,
2020 varentry2->ctx);
2021 if (r == -1)
2022 {
2023 /* ctx1 is a strict subset of ctx2, ignore ctx1. */
2024 matches[i] = false;
2025 break;
2026 }
2027 else if (r == 1)
2028 /* ctx2 is a strict subset of ctx1, remove ctx2. */
2029 matches[j] = false;
2030 }
2031 }
2032
2033 widest_int max_score = -1;
2034 varentry2 = NULL;
2035 FOR_EACH_VEC_SAFE_ELT (entryp->variants, i, varentry1)
2036 if (matches[i])
2037 {
2038 widest_int score
2039 = (cur_node->simdclone ? varentry1->score_in_declare_simd_clone
2040 : varentry1->score);
2041 if (score > max_score)
2042 {
2043 max_score = score;
2044 varentry2 = varentry1;
2045 }
2046 }
2047 return varentry2->variant->decl;
2048 }
2049
2050 /* Hook to adjust hash tables on cgraph_node removal. */
2051
2052 static void
2053 omp_declare_variant_remove_hook (struct cgraph_node *node, void *)
2054 {
2055 if (!node->declare_variant_alt)
2056 return;
2057
2058 /* Drop this hash table completely. */
2059 omp_declare_variants = NULL;
2060 /* And remove node from the other hash table. */
2061 if (omp_declare_variant_alt)
2062 {
2063 omp_declare_variant_base_entry entry;
2064 entry.base = NULL;
2065 entry.node = node;
2066 entry.variants = NULL;
2067 omp_declare_variant_alt->remove_elt_with_hash (&entry,
2068 DECL_UID (node->decl));
2069 }
2070 }
2071
2072 /* Try to resolve declare variant, return the variant decl if it should
2073 be used instead of base, or base otherwise. */
2074
2075 tree
2076 omp_resolve_declare_variant (tree base)
2077 {
2078 tree variant1 = NULL_TREE, variant2 = NULL_TREE;
2079 if (cfun && (cfun->curr_properties & PROP_gimple_any) != 0)
2080 return omp_resolve_late_declare_variant (base);
2081
2082 auto_vec <tree, 16> variants;
2083 auto_vec <bool, 16> defer;
2084 bool any_deferred = false;
2085 for (tree attr = DECL_ATTRIBUTES (base); attr; attr = TREE_CHAIN (attr))
2086 {
2087 attr = lookup_attribute ("omp declare variant base", attr);
2088 if (attr == NULL_TREE)
2089 break;
2090 if (TREE_CODE (TREE_PURPOSE (TREE_VALUE (attr))) != FUNCTION_DECL)
2091 continue;
2092 cgraph_node *node = cgraph_node::get (base);
2093 /* If this is already a magic decl created by this function,
2094 don't process it again. */
2095 if (node && node->declare_variant_alt)
2096 return base;
2097 switch (omp_context_selector_matches (TREE_VALUE (TREE_VALUE (attr))))
2098 {
2099 case 0:
2100 /* No match, ignore. */
2101 break;
2102 case -1:
2103 /* Needs to be deferred. */
2104 any_deferred = true;
2105 variants.safe_push (attr);
2106 defer.safe_push (true);
2107 break;
2108 default:
2109 variants.safe_push (attr);
2110 defer.safe_push (false);
2111 break;
2112 }
2113 }
2114 if (variants.length () == 0)
2115 return base;
2116
2117 if (any_deferred)
2118 {
2119 widest_int max_score1 = 0;
2120 widest_int max_score2 = 0;
2121 bool first = true;
2122 unsigned int i;
2123 tree attr1, attr2;
2124 omp_declare_variant_base_entry entry;
2125 entry.base = cgraph_node::get_create (base);
2126 entry.node = NULL;
2127 vec_alloc (entry.variants, variants.length ());
2128 FOR_EACH_VEC_ELT (variants, i, attr1)
2129 {
2130 widest_int score1;
2131 widest_int score2;
2132 bool need_two;
2133 tree ctx = TREE_VALUE (TREE_VALUE (attr1));
2134 need_two = omp_context_compute_score (ctx, &score1, false);
2135 if (need_two)
2136 omp_context_compute_score (ctx, &score2, true);
2137 else
2138 score2 = score1;
2139 if (first)
2140 {
2141 first = false;
2142 max_score1 = score1;
2143 max_score2 = score2;
2144 if (!defer[i])
2145 {
2146 variant1 = attr1;
2147 variant2 = attr1;
2148 }
2149 }
2150 else
2151 {
2152 if (max_score1 == score1)
2153 variant1 = NULL_TREE;
2154 else if (score1 > max_score1)
2155 {
2156 max_score1 = score1;
2157 variant1 = defer[i] ? NULL_TREE : attr1;
2158 }
2159 if (max_score2 == score2)
2160 variant2 = NULL_TREE;
2161 else if (score2 > max_score2)
2162 {
2163 max_score2 = score2;
2164 variant2 = defer[i] ? NULL_TREE : attr1;
2165 }
2166 }
2167 omp_declare_variant_entry varentry;
2168 varentry.variant
2169 = cgraph_node::get_create (TREE_PURPOSE (TREE_VALUE (attr1)));
2170 varentry.score = score1;
2171 varentry.score_in_declare_simd_clone = score2;
2172 varentry.ctx = ctx;
2173 varentry.matches = !defer[i];
2174 entry.variants->quick_push (varentry);
2175 }
2176
2177 /* If there is a clear winner variant with the score which is not
2178 deferred, verify it is not a strict subset of any other context
2179 selector and if it is not, it is the best alternative no matter
2180 whether the others do or don't match. */
2181 if (variant1 && variant1 == variant2)
2182 {
2183 tree ctx1 = TREE_VALUE (TREE_VALUE (variant1));
2184 FOR_EACH_VEC_ELT (variants, i, attr2)
2185 {
2186 if (attr2 == variant1)
2187 continue;
2188 tree ctx2 = TREE_VALUE (TREE_VALUE (attr2));
2189 int r = omp_context_selector_compare (ctx1, ctx2);
2190 if (r == -1)
2191 {
2192 /* The winner is a strict subset of ctx2, can't
2193 decide now. */
2194 variant1 = NULL_TREE;
2195 break;
2196 }
2197 }
2198 if (variant1)
2199 {
2200 vec_free (entry.variants);
2201 return TREE_PURPOSE (TREE_VALUE (variant1));
2202 }
2203 }
2204
2205 static struct cgraph_node_hook_list *node_removal_hook_holder;
2206 if (!node_removal_hook_holder)
2207 node_removal_hook_holder
2208 = symtab->add_cgraph_removal_hook (omp_declare_variant_remove_hook,
2209 NULL);
2210
2211 if (omp_declare_variants == NULL)
2212 omp_declare_variants
2213 = hash_table<omp_declare_variant_hasher>::create_ggc (64);
2214 omp_declare_variant_base_entry **slot
2215 = omp_declare_variants->find_slot (&entry, INSERT);
2216 if (*slot != NULL)
2217 {
2218 vec_free (entry.variants);
2219 return (*slot)->node->decl;
2220 }
2221
2222 *slot = ggc_cleared_alloc<omp_declare_variant_base_entry> ();
2223 (*slot)->base = entry.base;
2224 (*slot)->node = entry.base;
2225 (*slot)->variants = entry.variants;
2226 tree alt = build_decl (DECL_SOURCE_LOCATION (base), FUNCTION_DECL,
2227 DECL_NAME (base), TREE_TYPE (base));
2228 DECL_ARTIFICIAL (alt) = 1;
2229 DECL_IGNORED_P (alt) = 1;
2230 TREE_STATIC (alt) = 1;
2231 tree attributes = DECL_ATTRIBUTES (base);
2232 if (lookup_attribute ("noipa", attributes) == NULL)
2233 {
2234 attributes = tree_cons (get_identifier ("noipa"), NULL, attributes);
2235 if (lookup_attribute ("noinline", attributes) == NULL)
2236 attributes = tree_cons (get_identifier ("noinline"), NULL,
2237 attributes);
2238 if (lookup_attribute ("noclone", attributes) == NULL)
2239 attributes = tree_cons (get_identifier ("noclone"), NULL,
2240 attributes);
2241 if (lookup_attribute ("no_icf", attributes) == NULL)
2242 attributes = tree_cons (get_identifier ("no_icf"), NULL,
2243 attributes);
2244 }
2245 DECL_ATTRIBUTES (alt) = attributes;
2246 DECL_INITIAL (alt) = error_mark_node;
2247 (*slot)->node = cgraph_node::create (alt);
2248 (*slot)->node->declare_variant_alt = 1;
2249 (*slot)->node->create_reference (entry.base, IPA_REF_ADDR);
2250 omp_declare_variant_entry *varentry;
2251 FOR_EACH_VEC_SAFE_ELT (entry.variants, i, varentry)
2252 (*slot)->node->create_reference (varentry->variant, IPA_REF_ADDR);
2253 if (omp_declare_variant_alt == NULL)
2254 omp_declare_variant_alt
2255 = hash_table<omp_declare_variant_alt_hasher>::create_ggc (64);
2256 *omp_declare_variant_alt->find_slot_with_hash (*slot, DECL_UID (alt),
2257 INSERT) = *slot;
2258 return alt;
2259 }
2260
2261 if (variants.length () == 1)
2262 return TREE_PURPOSE (TREE_VALUE (variants[0]));
2263
2264 /* A context selector that is a strict subset of another context selector
2265 has a score of zero. */
2266 tree attr1, attr2;
2267 unsigned int i, j;
2268 FOR_EACH_VEC_ELT (variants, i, attr1)
2269 if (attr1)
2270 {
2271 tree ctx1 = TREE_VALUE (TREE_VALUE (attr1));
2272 FOR_EACH_VEC_ELT_FROM (variants, j, attr2, i + 1)
2273 if (attr2)
2274 {
2275 tree ctx2 = TREE_VALUE (TREE_VALUE (attr2));
2276 int r = omp_context_selector_compare (ctx1, ctx2);
2277 if (r == -1)
2278 {
2279 /* ctx1 is a strict subset of ctx2, remove
2280 attr1 from the vector. */
2281 variants[i] = NULL_TREE;
2282 break;
2283 }
2284 else if (r == 1)
2285 /* ctx2 is a strict subset of ctx1, remove attr2
2286 from the vector. */
2287 variants[j] = NULL_TREE;
2288 }
2289 }
2290 widest_int max_score1 = 0;
2291 widest_int max_score2 = 0;
2292 bool first = true;
2293 FOR_EACH_VEC_ELT (variants, i, attr1)
2294 if (attr1)
2295 {
2296 if (variant1)
2297 {
2298 widest_int score1;
2299 widest_int score2;
2300 bool need_two;
2301 tree ctx;
2302 if (first)
2303 {
2304 first = false;
2305 ctx = TREE_VALUE (TREE_VALUE (variant1));
2306 need_two = omp_context_compute_score (ctx, &max_score1, false);
2307 if (need_two)
2308 omp_context_compute_score (ctx, &max_score2, true);
2309 else
2310 max_score2 = max_score1;
2311 }
2312 ctx = TREE_VALUE (TREE_VALUE (attr1));
2313 need_two = omp_context_compute_score (ctx, &score1, false);
2314 if (need_two)
2315 omp_context_compute_score (ctx, &score2, true);
2316 else
2317 score2 = score1;
2318 if (score1 > max_score1)
2319 {
2320 max_score1 = score1;
2321 variant1 = attr1;
2322 }
2323 if (score2 > max_score2)
2324 {
2325 max_score2 = score2;
2326 variant2 = attr1;
2327 }
2328 }
2329 else
2330 {
2331 variant1 = attr1;
2332 variant2 = attr1;
2333 }
2334 }
2335 /* If there is a disagreement on which variant has the highest score
2336 depending on whether it will be in a declare simd clone or not,
2337 punt for now and defer until after IPA where we will know that. */
2338 return ((variant1 && variant1 == variant2)
2339 ? TREE_PURPOSE (TREE_VALUE (variant1)) : base);
2340 }
2341
2342 void
2343 omp_lto_output_declare_variant_alt (lto_simple_output_block *ob,
2344 cgraph_node *node,
2345 lto_symtab_encoder_t encoder)
2346 {
2347 gcc_assert (node->declare_variant_alt);
2348
2349 omp_declare_variant_base_entry entry;
2350 entry.base = NULL;
2351 entry.node = node;
2352 entry.variants = NULL;
2353 omp_declare_variant_base_entry *entryp
2354 = omp_declare_variant_alt->find_with_hash (&entry, DECL_UID (node->decl));
2355 gcc_assert (entryp);
2356
2357 int nbase = lto_symtab_encoder_lookup (encoder, entryp->base);
2358 gcc_assert (nbase != LCC_NOT_FOUND);
2359 streamer_write_hwi_stream (ob->main_stream, nbase);
2360
2361 streamer_write_hwi_stream (ob->main_stream, entryp->variants->length ());
2362
2363 unsigned int i;
2364 omp_declare_variant_entry *varentry;
2365 FOR_EACH_VEC_SAFE_ELT (entryp->variants, i, varentry)
2366 {
2367 int nvar = lto_symtab_encoder_lookup (encoder, varentry->variant);
2368 gcc_assert (nvar != LCC_NOT_FOUND);
2369 streamer_write_hwi_stream (ob->main_stream, nvar);
2370
2371 for (widest_int *w = &varentry->score; ;
2372 w = &varentry->score_in_declare_simd_clone)
2373 {
2374 unsigned len = w->get_len ();
2375 streamer_write_hwi_stream (ob->main_stream, len);
2376 const HOST_WIDE_INT *val = w->get_val ();
2377 for (unsigned j = 0; j < len; j++)
2378 streamer_write_hwi_stream (ob->main_stream, val[j]);
2379 if (w == &varentry->score_in_declare_simd_clone)
2380 break;
2381 }
2382
2383 HOST_WIDE_INT cnt = -1;
2384 HOST_WIDE_INT i = varentry->matches ? 1 : 0;
2385 for (tree attr = DECL_ATTRIBUTES (entryp->base->decl);
2386 attr; attr = TREE_CHAIN (attr), i += 2)
2387 {
2388 attr = lookup_attribute ("omp declare variant base", attr);
2389 if (attr == NULL_TREE)
2390 break;
2391
2392 if (varentry->ctx == TREE_VALUE (TREE_VALUE (attr)))
2393 {
2394 cnt = i;
2395 break;
2396 }
2397 }
2398
2399 gcc_assert (cnt != -1);
2400 streamer_write_hwi_stream (ob->main_stream, cnt);
2401 }
2402 }
2403
2404 void
2405 omp_lto_input_declare_variant_alt (lto_input_block *ib, cgraph_node *node,
2406 vec<symtab_node *> nodes)
2407 {
2408 gcc_assert (node->declare_variant_alt);
2409 omp_declare_variant_base_entry *entryp
2410 = ggc_cleared_alloc<omp_declare_variant_base_entry> ();
2411 entryp->base = dyn_cast<cgraph_node *> (nodes[streamer_read_hwi (ib)]);
2412 entryp->node = node;
2413 unsigned int len = streamer_read_hwi (ib);
2414 vec_alloc (entryp->variants, len);
2415
2416 for (unsigned int i = 0; i < len; i++)
2417 {
2418 omp_declare_variant_entry varentry;
2419 varentry.variant
2420 = dyn_cast<cgraph_node *> (nodes[streamer_read_hwi (ib)]);
2421 for (widest_int *w = &varentry.score; ;
2422 w = &varentry.score_in_declare_simd_clone)
2423 {
2424 unsigned len2 = streamer_read_hwi (ib);
2425 HOST_WIDE_INT arr[WIDE_INT_MAX_ELTS];
2426 gcc_assert (len2 <= WIDE_INT_MAX_ELTS);
2427 for (unsigned int j = 0; j < len2; j++)
2428 arr[j] = streamer_read_hwi (ib);
2429 *w = widest_int::from_array (arr, len2, true);
2430 if (w == &varentry.score_in_declare_simd_clone)
2431 break;
2432 }
2433
2434 HOST_WIDE_INT cnt = streamer_read_hwi (ib);
2435 HOST_WIDE_INT j = 0;
2436 varentry.ctx = NULL_TREE;
2437 varentry.matches = (cnt & 1) ? true : false;
2438 cnt &= ~HOST_WIDE_INT_1;
2439 for (tree attr = DECL_ATTRIBUTES (entryp->base->decl);
2440 attr; attr = TREE_CHAIN (attr), j += 2)
2441 {
2442 attr = lookup_attribute ("omp declare variant base", attr);
2443 if (attr == NULL_TREE)
2444 break;
2445
2446 if (cnt == j)
2447 {
2448 varentry.ctx = TREE_VALUE (TREE_VALUE (attr));
2449 break;
2450 }
2451 }
2452 gcc_assert (varentry.ctx != NULL_TREE);
2453 entryp->variants->quick_push (varentry);
2454 }
2455 if (omp_declare_variant_alt == NULL)
2456 omp_declare_variant_alt
2457 = hash_table<omp_declare_variant_alt_hasher>::create_ggc (64);
2458 *omp_declare_variant_alt->find_slot_with_hash (entryp, DECL_UID (node->decl),
2459 INSERT) = entryp;
2460 }
2461
2462 /* Encode an oacc launch argument. This matches the GOMP_LAUNCH_PACK
2463 macro on gomp-constants.h. We do not check for overflow. */
2464
2465 tree
2466 oacc_launch_pack (unsigned code, tree device, unsigned op)
2467 {
2468 tree res;
2469
2470 res = build_int_cst (unsigned_type_node, GOMP_LAUNCH_PACK (code, 0, op));
2471 if (device)
2472 {
2473 device = fold_build2 (LSHIFT_EXPR, unsigned_type_node,
2474 device, build_int_cst (unsigned_type_node,
2475 GOMP_LAUNCH_DEVICE_SHIFT));
2476 res = fold_build2 (BIT_IOR_EXPR, unsigned_type_node, res, device);
2477 }
2478 return res;
2479 }
2480
2481 /* FIXME: What is the following comment for? */
2482 /* Look for compute grid dimension clauses and convert to an attribute
2483 attached to FN. This permits the target-side code to (a) massage
2484 the dimensions, (b) emit that data and (c) optimize. Non-constant
2485 dimensions are pushed onto ARGS.
2486
2487 The attribute value is a TREE_LIST. A set of dimensions is
2488 represented as a list of INTEGER_CST. Those that are runtime
2489 exprs are represented as an INTEGER_CST of zero.
2490
2491 TODO: Normally the attribute will just contain a single such list. If
2492 however it contains a list of lists, this will represent the use of
2493 device_type. Each member of the outer list is an assoc list of
2494 dimensions, keyed by the device type. The first entry will be the
2495 default. Well, that's the plan. */
2496
2497 /* Replace any existing oacc fn attribute with updated dimensions. */
2498
2499 /* Variant working on a list of attributes. */
2500
2501 tree
2502 oacc_replace_fn_attrib_attr (tree attribs, tree dims)
2503 {
2504 tree ident = get_identifier (OACC_FN_ATTRIB);
2505
2506 /* If we happen to be present as the first attrib, drop it. */
2507 if (attribs && TREE_PURPOSE (attribs) == ident)
2508 attribs = TREE_CHAIN (attribs);
2509 return tree_cons (ident, dims, attribs);
2510 }
2511
2512 /* Variant working on a function decl. */
2513
2514 void
2515 oacc_replace_fn_attrib (tree fn, tree dims)
2516 {
2517 DECL_ATTRIBUTES (fn)
2518 = oacc_replace_fn_attrib_attr (DECL_ATTRIBUTES (fn), dims);
2519 }
2520
2521 /* Scan CLAUSES for launch dimensions and attach them to the oacc
2522 function attribute. Push any that are non-constant onto the ARGS
2523 list, along with an appropriate GOMP_LAUNCH_DIM tag. */
2524
2525 void
2526 oacc_set_fn_attrib (tree fn, tree clauses, vec<tree> *args)
2527 {
2528 /* Must match GOMP_DIM ordering. */
2529 static const omp_clause_code ids[]
2530 = { OMP_CLAUSE_NUM_GANGS, OMP_CLAUSE_NUM_WORKERS,
2531 OMP_CLAUSE_VECTOR_LENGTH };
2532 unsigned ix;
2533 tree dims[GOMP_DIM_MAX];
2534
2535 tree attr = NULL_TREE;
2536 unsigned non_const = 0;
2537
2538 for (ix = GOMP_DIM_MAX; ix--;)
2539 {
2540 tree clause = omp_find_clause (clauses, ids[ix]);
2541 tree dim = NULL_TREE;
2542
2543 if (clause)
2544 dim = OMP_CLAUSE_EXPR (clause, ids[ix]);
2545 dims[ix] = dim;
2546 if (dim && TREE_CODE (dim) != INTEGER_CST)
2547 {
2548 dim = integer_zero_node;
2549 non_const |= GOMP_DIM_MASK (ix);
2550 }
2551 attr = tree_cons (NULL_TREE, dim, attr);
2552 }
2553
2554 oacc_replace_fn_attrib (fn, attr);
2555
2556 if (non_const)
2557 {
2558 /* Push a dynamic argument set. */
2559 args->safe_push (oacc_launch_pack (GOMP_LAUNCH_DIM,
2560 NULL_TREE, non_const));
2561 for (unsigned ix = 0; ix != GOMP_DIM_MAX; ix++)
2562 if (non_const & GOMP_DIM_MASK (ix))
2563 args->safe_push (dims[ix]);
2564 }
2565 }
2566
2567 /* Verify OpenACC routine clauses.
2568
2569 Returns 0 if FNDECL should be marked with an OpenACC 'routine' directive, 1
2570 if it has already been marked in compatible way, and -1 if incompatible.
2571 Upon returning, the chain of clauses will contain exactly one clause
2572 specifying the level of parallelism. */
2573
2574 int
2575 oacc_verify_routine_clauses (tree fndecl, tree *clauses, location_t loc,
2576 const char *routine_str)
2577 {
2578 tree c_level = NULL_TREE;
2579 tree c_p = NULL_TREE;
2580 for (tree c = *clauses; c; c_p = c, c = OMP_CLAUSE_CHAIN (c))
2581 switch (OMP_CLAUSE_CODE (c))
2582 {
2583 case OMP_CLAUSE_GANG:
2584 case OMP_CLAUSE_WORKER:
2585 case OMP_CLAUSE_VECTOR:
2586 case OMP_CLAUSE_SEQ:
2587 if (c_level == NULL_TREE)
2588 c_level = c;
2589 else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_CODE (c_level))
2590 {
2591 /* This has already been diagnosed in the front ends. */
2592 /* Drop the duplicate clause. */
2593 gcc_checking_assert (c_p != NULL_TREE);
2594 OMP_CLAUSE_CHAIN (c_p) = OMP_CLAUSE_CHAIN (c);
2595 c = c_p;
2596 }
2597 else
2598 {
2599 error_at (OMP_CLAUSE_LOCATION (c),
2600 "%qs specifies a conflicting level of parallelism",
2601 omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
2602 inform (OMP_CLAUSE_LOCATION (c_level),
2603 "... to the previous %qs clause here",
2604 omp_clause_code_name[OMP_CLAUSE_CODE (c_level)]);
2605 /* Drop the conflicting clause. */
2606 gcc_checking_assert (c_p != NULL_TREE);
2607 OMP_CLAUSE_CHAIN (c_p) = OMP_CLAUSE_CHAIN (c);
2608 c = c_p;
2609 }
2610 break;
2611 default:
2612 gcc_unreachable ();
2613 }
2614 if (c_level == NULL_TREE)
2615 {
2616 /* Default to an implicit 'seq' clause. */
2617 c_level = build_omp_clause (loc, OMP_CLAUSE_SEQ);
2618 OMP_CLAUSE_CHAIN (c_level) = *clauses;
2619 *clauses = c_level;
2620 }
2621 /* In *clauses, we now have exactly one clause specifying the level of
2622 parallelism. */
2623
2624 tree attr
2625 = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (fndecl));
2626 if (attr != NULL_TREE)
2627 {
2628 /* Diagnose if "#pragma omp declare target" has also been applied. */
2629 if (TREE_VALUE (attr) == NULL_TREE)
2630 {
2631 /* See <https://gcc.gnu.org/PR93465>; the semantics of combining
2632 OpenACC and OpenMP 'target' are not clear. */
2633 error_at (loc,
2634 "cannot apply %<%s%> to %qD, which has also been"
2635 " marked with an OpenMP 'declare target' directive",
2636 routine_str, fndecl);
2637 /* Incompatible. */
2638 return -1;
2639 }
2640
2641 /* If a "#pragma acc routine" has already been applied, just verify
2642 this one for compatibility. */
2643 /* Collect previous directive's clauses. */
2644 tree c_level_p = NULL_TREE;
2645 for (tree c = TREE_VALUE (attr); c; c = OMP_CLAUSE_CHAIN (c))
2646 switch (OMP_CLAUSE_CODE (c))
2647 {
2648 case OMP_CLAUSE_GANG:
2649 case OMP_CLAUSE_WORKER:
2650 case OMP_CLAUSE_VECTOR:
2651 case OMP_CLAUSE_SEQ:
2652 gcc_checking_assert (c_level_p == NULL_TREE);
2653 c_level_p = c;
2654 break;
2655 default:
2656 gcc_unreachable ();
2657 }
2658 gcc_checking_assert (c_level_p != NULL_TREE);
2659 /* ..., and compare to current directive's, which we've already collected
2660 above. */
2661 tree c_diag;
2662 tree c_diag_p;
2663 /* Matching level of parallelism? */
2664 if (OMP_CLAUSE_CODE (c_level) != OMP_CLAUSE_CODE (c_level_p))
2665 {
2666 c_diag = c_level;
2667 c_diag_p = c_level_p;
2668 goto incompatible;
2669 }
2670 /* Compatible. */
2671 return 1;
2672
2673 incompatible:
2674 if (c_diag != NULL_TREE)
2675 error_at (OMP_CLAUSE_LOCATION (c_diag),
2676 "incompatible %qs clause when applying"
2677 " %<%s%> to %qD, which has already been"
2678 " marked with an OpenACC 'routine' directive",
2679 omp_clause_code_name[OMP_CLAUSE_CODE (c_diag)],
2680 routine_str, fndecl);
2681 else if (c_diag_p != NULL_TREE)
2682 error_at (loc,
2683 "missing %qs clause when applying"
2684 " %<%s%> to %qD, which has already been"
2685 " marked with an OpenACC 'routine' directive",
2686 omp_clause_code_name[OMP_CLAUSE_CODE (c_diag_p)],
2687 routine_str, fndecl);
2688 else
2689 gcc_unreachable ();
2690 if (c_diag_p != NULL_TREE)
2691 inform (OMP_CLAUSE_LOCATION (c_diag_p),
2692 "... with %qs clause here",
2693 omp_clause_code_name[OMP_CLAUSE_CODE (c_diag_p)]);
2694 else
2695 {
2696 /* In the front ends, we don't preserve location information for the
2697 OpenACC routine directive itself. However, that of c_level_p
2698 should be close. */
2699 location_t loc_routine = OMP_CLAUSE_LOCATION (c_level_p);
2700 inform (loc_routine, "... without %qs clause near to here",
2701 omp_clause_code_name[OMP_CLAUSE_CODE (c_diag)]);
2702 }
2703 /* Incompatible. */
2704 return -1;
2705 }
2706
2707 return 0;
2708 }
2709
2710 /* Process the OpenACC 'routine' directive clauses to generate an attribute
2711 for the level of parallelism. All dimensions have a size of zero
2712 (dynamic). TREE_PURPOSE is set to indicate whether that dimension
2713 can have a loop partitioned on it. non-zero indicates
2714 yes, zero indicates no. By construction once a non-zero has been
2715 reached, further inner dimensions must also be non-zero. We set
2716 TREE_VALUE to zero for the dimensions that may be partitioned and
2717 1 for the other ones -- if a loop is (erroneously) spawned at
2718 an outer level, we don't want to try and partition it. */
2719
2720 tree
2721 oacc_build_routine_dims (tree clauses)
2722 {
2723 /* Must match GOMP_DIM ordering. */
2724 static const omp_clause_code ids[]
2725 = {OMP_CLAUSE_GANG, OMP_CLAUSE_WORKER, OMP_CLAUSE_VECTOR, OMP_CLAUSE_SEQ};
2726 int ix;
2727 int level = -1;
2728
2729 for (; clauses; clauses = OMP_CLAUSE_CHAIN (clauses))
2730 for (ix = GOMP_DIM_MAX + 1; ix--;)
2731 if (OMP_CLAUSE_CODE (clauses) == ids[ix])
2732 {
2733 level = ix;
2734 break;
2735 }
2736 gcc_checking_assert (level >= 0);
2737
2738 tree dims = NULL_TREE;
2739
2740 for (ix = GOMP_DIM_MAX; ix--;)
2741 dims = tree_cons (build_int_cst (boolean_type_node, ix >= level),
2742 build_int_cst (integer_type_node, ix < level), dims);
2743
2744 return dims;
2745 }
2746
2747 /* Retrieve the oacc function attrib and return it. Non-oacc
2748 functions will return NULL. */
2749
2750 tree
2751 oacc_get_fn_attrib (tree fn)
2752 {
2753 return lookup_attribute (OACC_FN_ATTRIB, DECL_ATTRIBUTES (fn));
2754 }
2755
2756 /* Return true if FN is an OpenMP or OpenACC offloading function. */
2757
2758 bool
2759 offloading_function_p (tree fn)
2760 {
2761 tree attrs = DECL_ATTRIBUTES (fn);
2762 return (lookup_attribute ("omp declare target", attrs)
2763 || lookup_attribute ("omp target entrypoint", attrs));
2764 }
2765
2766 /* Extract an oacc execution dimension from FN. FN must be an
2767 offloaded function or routine that has already had its execution
2768 dimensions lowered to the target-specific values. */
2769
2770 int
2771 oacc_get_fn_dim_size (tree fn, int axis)
2772 {
2773 tree attrs = oacc_get_fn_attrib (fn);
2774
2775 gcc_assert (axis < GOMP_DIM_MAX);
2776
2777 tree dims = TREE_VALUE (attrs);
2778 while (axis--)
2779 dims = TREE_CHAIN (dims);
2780
2781 int size = TREE_INT_CST_LOW (TREE_VALUE (dims));
2782
2783 return size;
2784 }
2785
2786 /* Extract the dimension axis from an IFN_GOACC_DIM_POS or
2787 IFN_GOACC_DIM_SIZE call. */
2788
2789 int
2790 oacc_get_ifn_dim_arg (const gimple *stmt)
2791 {
2792 gcc_checking_assert (gimple_call_internal_fn (stmt) == IFN_GOACC_DIM_SIZE
2793 || gimple_call_internal_fn (stmt) == IFN_GOACC_DIM_POS);
2794 tree arg = gimple_call_arg (stmt, 0);
2795 HOST_WIDE_INT axis = TREE_INT_CST_LOW (arg);
2796
2797 gcc_checking_assert (axis >= 0 && axis < GOMP_DIM_MAX);
2798 return (int) axis;
2799 }
2800
2801 #include "gt-omp-general.h"
Mirror of the offical gcc git repository hosted at gcc.gnu.org