LCOV - code coverage report
Current view: top level - gcc - omp-general.c (source / functions) Hit Total Coverage
Test: gcc.info Lines: 866 930 93.1 %
Date: 2020-04-04 11:58:09 Functions: 30 32 93.8 %
Legend: Lines: hit not hit | Branches: + taken - not taken # not executed Branches: 0 0 -

           Branch data     Line data    Source code
       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 "hsa-common.h"
      43                 :            : #include "tree-pass.h"
      44                 :            : #include "omp-device-properties.h"
      45                 :            : 
      46                 :            : enum omp_requires omp_requires_mask;
      47                 :            : 
      48                 :            : tree
      49                 :     772247 : omp_find_clause (tree clauses, enum omp_clause_code kind)
      50                 :            : {
      51                 :    2728820 :   for (; clauses ; clauses = OMP_CLAUSE_CHAIN (clauses))
      52                 :    2135960 :     if (OMP_CLAUSE_CODE (clauses) == kind)
      53                 :     179392 :       return clauses;
      54                 :            : 
      55                 :            :   return NULL_TREE;
      56                 :            : }
      57                 :            : 
      58                 :            : /* True if OpenMP should regard this DECL as being a scalar which has Fortran's
      59                 :            :    allocatable or pointer attribute.  */
      60                 :            : bool
      61                 :      11564 : omp_is_allocatable_or_ptr (tree decl)
      62                 :            : {
      63                 :      11564 :   return lang_hooks.decls.omp_is_allocatable_or_ptr (decl);
      64                 :            : }
      65                 :            : 
      66                 :            : /* Check whether this DECL belongs to a Fortran optional argument.
      67                 :            :    With 'for_present_check' set to false, decls which are optional parameters
      68                 :            :    themselve are returned as tree - or a NULL_TREE otherwise. Those decls are
      69                 :            :    always pointers.  With 'for_present_check' set to true, the decl for checking
      70                 :            :    whether an argument is present is returned; for arguments with value
      71                 :            :    attribute this is the hidden argument and of BOOLEAN_TYPE.  If the decl is
      72                 :            :    unrelated to optional arguments, NULL_TREE is returned.  */
      73                 :            : 
      74                 :            : tree
      75                 :      10384 : omp_check_optional_argument (tree decl, bool for_present_check)
      76                 :            : {
      77                 :      10384 :   return lang_hooks.decls.omp_check_optional_argument (decl, for_present_check);
      78                 :            : }
      79                 :            : 
      80                 :            : /* Return true if DECL is a reference type.  */
      81                 :            : 
      82                 :            : bool
      83                 :     572730 : omp_is_reference (tree decl)
      84                 :            : {
      85                 :     572730 :   return lang_hooks.decls.omp_privatize_by_reference (decl);
      86                 :            : }
      87                 :            : 
      88                 :            : /* Adjust *COND_CODE and *N2 so that the former is either LT_EXPR or GT_EXPR,
      89                 :            :    given that V is the loop index variable and STEP is loop step. */
      90                 :            : 
      91                 :            : void
      92                 :     156148 : omp_adjust_for_condition (location_t loc, enum tree_code *cond_code, tree *n2,
      93                 :            :                           tree v, tree step)
      94                 :            : {
      95                 :     156148 :   switch (*cond_code)
      96                 :            :     {
      97                 :            :     case LT_EXPR:
      98                 :            :     case GT_EXPR:
      99                 :            :       break;
     100                 :            : 
     101                 :      29237 :     case NE_EXPR:
     102                 :      29237 :       gcc_assert (TREE_CODE (step) == INTEGER_CST);
     103                 :      29237 :       if (TREE_CODE (TREE_TYPE (v)) == INTEGER_TYPE)
     104                 :            :         {
     105                 :      24530 :           if (integer_onep (step))
     106                 :      17240 :             *cond_code = LT_EXPR;
     107                 :            :           else
     108                 :            :             {
     109                 :       7290 :               gcc_assert (integer_minus_onep (step));
     110                 :       7290 :               *cond_code = GT_EXPR;
     111                 :            :             }
     112                 :            :         }
     113                 :            :       else
     114                 :            :         {
     115                 :       4707 :           tree unit = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (v)));
     116                 :       4707 :           gcc_assert (TREE_CODE (unit) == INTEGER_CST);
     117                 :       4707 :           if (tree_int_cst_equal (unit, step))
     118                 :       3221 :             *cond_code = LT_EXPR;
     119                 :            :           else
     120                 :            :             {
     121                 :       1486 :               gcc_assert (wi::neg (wi::to_widest (unit))
     122                 :            :                           == wi::to_widest (step));
     123                 :       1486 :               *cond_code = GT_EXPR;
     124                 :            :             }
     125                 :            :         }
     126                 :            : 
     127                 :            :       break;
     128                 :            : 
     129                 :      15506 :     case LE_EXPR:
     130                 :      15506 :       if (POINTER_TYPE_P (TREE_TYPE (*n2)))
     131                 :        102 :         *n2 = fold_build_pointer_plus_hwi_loc (loc, *n2, 1);
     132                 :            :       else
     133                 :      15404 :         *n2 = fold_build2_loc (loc, PLUS_EXPR, TREE_TYPE (*n2), *n2,
     134                 :      15404 :                                build_int_cst (TREE_TYPE (*n2), 1));
     135                 :      15506 :       *cond_code = LT_EXPR;
     136                 :      15506 :       break;
     137                 :        675 :     case GE_EXPR:
     138                 :        675 :       if (POINTER_TYPE_P (TREE_TYPE (*n2)))
     139                 :         96 :         *n2 = fold_build_pointer_plus_hwi_loc (loc, *n2, -1);
     140                 :            :       else
     141                 :        579 :         *n2 = fold_build2_loc (loc, MINUS_EXPR, TREE_TYPE (*n2), *n2,
     142                 :        579 :                                build_int_cst (TREE_TYPE (*n2), 1));
     143                 :        675 :       *cond_code = GT_EXPR;
     144                 :        675 :       break;
     145                 :          0 :     default:
     146                 :          0 :       gcc_unreachable ();
     147                 :            :     }
     148                 :     156148 : }
     149                 :            : 
     150                 :            : /* Return the looping step from INCR, extracted from the step of a gimple omp
     151                 :            :    for statement.  */
     152                 :            : 
     153                 :            : tree
     154                 :     156148 : omp_get_for_step_from_incr (location_t loc, tree incr)
     155                 :            : {
     156                 :     156148 :   tree step;
     157                 :     156148 :   switch (TREE_CODE (incr))
     158                 :            :     {
     159                 :     134367 :     case PLUS_EXPR:
     160                 :     134367 :       step = TREE_OPERAND (incr, 1);
     161                 :     134367 :       break;
     162                 :      15763 :     case POINTER_PLUS_EXPR:
     163                 :      15763 :       step = fold_convert (ssizetype, TREE_OPERAND (incr, 1));
     164                 :      15763 :       break;
     165                 :       6018 :     case MINUS_EXPR:
     166                 :       6018 :       step = TREE_OPERAND (incr, 1);
     167                 :       6018 :       step = fold_build1_loc (loc, NEGATE_EXPR, TREE_TYPE (step), step);
     168                 :       6018 :       break;
     169                 :          0 :     default:
     170                 :          0 :       gcc_unreachable ();
     171                 :            :     }
     172                 :     156148 :   return step;
     173                 :            : }
     174                 :            : 
     175                 :            : /* Extract the header elements of parallel loop FOR_STMT and store
     176                 :            :    them into *FD.  */
     177                 :            : 
     178                 :            : void
     179                 :      99230 : omp_extract_for_data (gomp_for *for_stmt, struct omp_for_data *fd,
     180                 :            :                       struct omp_for_data_loop *loops)
     181                 :            : {
     182                 :      99230 :   tree t, var, *collapse_iter, *collapse_count;
     183                 :      99230 :   tree count = NULL_TREE, iter_type = long_integer_type_node;
     184                 :      99230 :   struct omp_for_data_loop *loop;
     185                 :      99230 :   int i;
     186                 :      99230 :   struct omp_for_data_loop dummy_loop;
     187                 :      99230 :   location_t loc = gimple_location (for_stmt);
     188                 :      99230 :   bool simd = gimple_omp_for_kind (for_stmt) == GF_OMP_FOR_KIND_SIMD;
     189                 :      99230 :   bool distribute = gimple_omp_for_kind (for_stmt)
     190                 :      99230 :                     == GF_OMP_FOR_KIND_DISTRIBUTE;
     191                 :      99230 :   bool taskloop = gimple_omp_for_kind (for_stmt)
     192                 :      99230 :                   == GF_OMP_FOR_KIND_TASKLOOP;
     193                 :      99230 :   tree iterv, countv;
     194                 :            : 
     195                 :      99230 :   fd->for_stmt = for_stmt;
     196                 :      99230 :   fd->pre = NULL;
     197                 :      99230 :   fd->have_nowait = distribute || simd;
     198                 :      99230 :   fd->have_ordered = false;
     199                 :      99230 :   fd->have_reductemp = false;
     200                 :      99230 :   fd->have_pointer_condtemp = false;
     201                 :      99230 :   fd->have_scantemp = false;
     202                 :      99230 :   fd->have_nonctrl_scantemp = false;
     203                 :      99230 :   fd->lastprivate_conditional = 0;
     204                 :      99230 :   fd->tiling = NULL_TREE;
     205                 :      99230 :   fd->collapse = 1;
     206                 :      99230 :   fd->ordered = 0;
     207                 :      99230 :   fd->sched_kind = OMP_CLAUSE_SCHEDULE_STATIC;
     208                 :      99230 :   fd->sched_modifiers = 0;
     209                 :      99230 :   fd->chunk_size = NULL_TREE;
     210                 :      99230 :   fd->simd_schedule = false;
     211                 :      99230 :   collapse_iter = NULL;
     212                 :      99230 :   collapse_count = NULL;
     213                 :            : 
     214                 :     508044 :   for (t = gimple_omp_for_clauses (for_stmt); t ; t = OMP_CLAUSE_CHAIN (t))
     215                 :     408814 :     switch (OMP_CLAUSE_CODE (t))
     216                 :            :       {
     217                 :      37568 :       case OMP_CLAUSE_NOWAIT:
     218                 :      37568 :         fd->have_nowait = true;
     219                 :      37568 :         break;
     220                 :       1092 :       case OMP_CLAUSE_ORDERED:
     221                 :       1092 :         fd->have_ordered = true;
     222                 :       1092 :         if (OMP_CLAUSE_ORDERED_EXPR (t))
     223                 :        533 :           fd->ordered = tree_to_shwi (OMP_CLAUSE_ORDERED_EXPR (t));
     224                 :            :         break;
     225                 :      29108 :       case OMP_CLAUSE_SCHEDULE:
     226                 :      29108 :         gcc_assert (!distribute && !taskloop);
     227                 :      29108 :         fd->sched_kind
     228                 :      29108 :           = (enum omp_clause_schedule_kind)
     229                 :      29108 :             (OMP_CLAUSE_SCHEDULE_KIND (t) & OMP_CLAUSE_SCHEDULE_MASK);
     230                 :      29108 :         fd->sched_modifiers = (OMP_CLAUSE_SCHEDULE_KIND (t)
     231                 :      29108 :                                & ~OMP_CLAUSE_SCHEDULE_MASK);
     232                 :      29108 :         fd->chunk_size = OMP_CLAUSE_SCHEDULE_CHUNK_EXPR (t);
     233                 :      29108 :         fd->simd_schedule = OMP_CLAUSE_SCHEDULE_SIMD (t);
     234                 :      29108 :         break;
     235                 :       6488 :       case OMP_CLAUSE_DIST_SCHEDULE:
     236                 :       6488 :         gcc_assert (distribute);
     237                 :       6488 :         fd->chunk_size = OMP_CLAUSE_DIST_SCHEDULE_CHUNK_EXPR (t);
     238                 :       6488 :         break;
     239                 :      30631 :       case OMP_CLAUSE_COLLAPSE:
     240                 :      30631 :         fd->collapse = tree_to_shwi (OMP_CLAUSE_COLLAPSE_EXPR (t));
     241                 :      30631 :         if (fd->collapse > 1)
     242                 :            :           {
     243                 :      29007 :             collapse_iter = &OMP_CLAUSE_COLLAPSE_ITERVAR (t);
     244                 :      29007 :             collapse_count = &OMP_CLAUSE_COLLAPSE_COUNT (t);
     245                 :            :           }
     246                 :            :         break;
     247                 :        354 :       case OMP_CLAUSE_TILE:
     248                 :        354 :         fd->tiling = OMP_CLAUSE_TILE_LIST (t);
     249                 :        354 :         fd->collapse = list_length (fd->tiling);
     250                 :        354 :         gcc_assert (fd->collapse);
     251                 :        354 :         collapse_iter = &OMP_CLAUSE_TILE_ITERVAR (t);
     252                 :        354 :         collapse_count = &OMP_CLAUSE_TILE_COUNT (t);
     253                 :        354 :         break;
     254                 :        370 :       case OMP_CLAUSE__REDUCTEMP_:
     255                 :        370 :         fd->have_reductemp = true;
     256                 :        370 :         break;
     257                 :      28633 :       case OMP_CLAUSE_LASTPRIVATE:
     258                 :      28633 :         if (OMP_CLAUSE_LASTPRIVATE_CONDITIONAL (t))
     259                 :       1210 :           fd->lastprivate_conditional++;
     260                 :            :         break;
     261                 :       1222 :       case OMP_CLAUSE__CONDTEMP_:
     262                 :       1222 :         if (POINTER_TYPE_P (TREE_TYPE (OMP_CLAUSE_DECL (t))))
     263                 :        182 :           fd->have_pointer_condtemp = true;
     264                 :            :         break;
     265                 :        994 :       case OMP_CLAUSE__SCANTEMP_:
     266                 :        994 :         fd->have_scantemp = true;
     267                 :        994 :         if (!OMP_CLAUSE__SCANTEMP__ALLOC (t)
     268                 :        994 :             && !OMP_CLAUSE__SCANTEMP__CONTROL (t))
     269                 :        187 :           fd->have_nonctrl_scantemp = true;
     270                 :            :         break;
     271                 :            :       default:
     272                 :            :         break;
     273                 :            :       }
     274                 :            : 
     275                 :      99230 :   if (fd->collapse > 1 || fd->tiling)
     276                 :      29356 :     fd->loops = loops;
     277                 :            :   else
     278                 :      69874 :     fd->loops = &fd->loop;
     279                 :            : 
     280                 :      99230 :   if (fd->ordered && fd->collapse == 1 && loops != NULL)
     281                 :            :     {
     282                 :        142 :       fd->loops = loops;
     283                 :        142 :       iterv = NULL_TREE;
     284                 :        142 :       countv = NULL_TREE;
     285                 :        142 :       collapse_iter = &iterv;
     286                 :        142 :       collapse_count = &countv;
     287                 :            :     }
     288                 :            : 
     289                 :            :   /* FIXME: for now map schedule(auto) to schedule(static).
     290                 :            :      There should be analysis to determine whether all iterations
     291                 :            :      are approximately the same amount of work (then schedule(static)
     292                 :            :      is best) or if it varies (then schedule(dynamic,N) is better).  */
     293                 :      99230 :   if (fd->sched_kind == OMP_CLAUSE_SCHEDULE_AUTO)
     294                 :            :     {
     295                 :       5722 :       fd->sched_kind = OMP_CLAUSE_SCHEDULE_STATIC;
     296                 :       5722 :       gcc_assert (fd->chunk_size == NULL);
     297                 :            :     }
     298                 :      99230 :   gcc_assert ((fd->collapse == 1 && !fd->tiling) || collapse_iter != NULL);
     299                 :      99230 :   if (taskloop)
     300                 :       4337 :     fd->sched_kind = OMP_CLAUSE_SCHEDULE_RUNTIME;
     301                 :      99230 :   if (fd->sched_kind == OMP_CLAUSE_SCHEDULE_RUNTIME)
     302                 :      11388 :     gcc_assert (fd->chunk_size == NULL);
     303                 :      87842 :   else if (fd->chunk_size == NULL)
     304                 :            :     {
     305                 :            :       /* We only need to compute a default chunk size for ordered
     306                 :            :          static loops and dynamic loops.  */
     307                 :      66393 :       if (fd->sched_kind != OMP_CLAUSE_SCHEDULE_STATIC
     308                 :      65869 :           || fd->have_ordered)
     309                 :       1036 :         fd->chunk_size = (fd->sched_kind == OMP_CLAUSE_SCHEDULE_STATIC)
     310                 :       1036 :                          ? integer_zero_node : integer_one_node;
     311                 :            :     }
     312                 :            : 
     313                 :      99230 :   int cnt = fd->ordered ? fd->ordered : fd->collapse;
     314                 :     255378 :   for (i = 0; i < cnt; i++)
     315                 :            :     {
     316                 :     156148 :       if (i == 0
     317                 :      99230 :           && fd->collapse == 1
     318                 :      70072 :           && !fd->tiling
     319                 :      69874 :           && (fd->ordered == 0 || loops == NULL))
     320                 :      69732 :         loop = &fd->loop;
     321                 :      86416 :       else if (loops != NULL)
     322                 :      28012 :         loop = loops + i;
     323                 :            :       else
     324                 :            :         loop = &dummy_loop;
     325                 :            : 
     326                 :     156148 :       loop->v = gimple_omp_for_index (for_stmt, i);
     327                 :     156148 :       gcc_assert (SSA_VAR_P (loop->v));
     328                 :     156148 :       gcc_assert (TREE_CODE (TREE_TYPE (loop->v)) == INTEGER_TYPE
     329                 :            :                   || TREE_CODE (TREE_TYPE (loop->v)) == POINTER_TYPE);
     330                 :     156148 :       var = TREE_CODE (loop->v) == SSA_NAME ? SSA_NAME_VAR (loop->v) : loop->v;
     331                 :     156148 :       loop->n1 = gimple_omp_for_initial (for_stmt, i);
     332                 :            : 
     333                 :     156148 :       loop->cond_code = gimple_omp_for_cond (for_stmt, i);
     334                 :     156148 :       loop->n2 = gimple_omp_for_final (for_stmt, i);
     335                 :     156148 :       gcc_assert (loop->cond_code != NE_EXPR
     336                 :            :                   || (gimple_omp_for_kind (for_stmt)
     337                 :            :                       != GF_OMP_FOR_KIND_OACC_LOOP));
     338                 :            : 
     339                 :     156148 :       t = gimple_omp_for_incr (for_stmt, i);
     340                 :     156148 :       gcc_assert (TREE_OPERAND (t, 0) == var);
     341                 :     156148 :       loop->step = omp_get_for_step_from_incr (loc, t);
     342                 :            : 
     343                 :     156148 :       omp_adjust_for_condition (loc, &loop->cond_code, &loop->n2, loop->v,
     344                 :            :                                 loop->step);
     345                 :            : 
     346                 :     156148 :       if (simd
     347                 :     125232 :           || (fd->sched_kind == OMP_CLAUSE_SCHEDULE_STATIC
     348                 :      92555 :               && !fd->have_ordered))
     349                 :            :         {
     350                 :     122256 :           if (fd->collapse == 1 && !fd->tiling)
     351                 :      56944 :             iter_type = TREE_TYPE (loop->v);
     352                 :      65312 :           else if (i == 0
     353                 :      65312 :                    || TYPE_PRECISION (iter_type)
     354                 :      42684 :                       < TYPE_PRECISION (TREE_TYPE (loop->v)))
     355                 :      33499 :             iter_type
     356                 :            :               = build_nonstandard_integer_type
     357                 :      33499 :                   (TYPE_PRECISION (TREE_TYPE (loop->v)), 1);
     358                 :            :         }
     359                 :      33892 :       else if (iter_type != long_long_unsigned_type_node)
     360                 :            :         {
     361                 :      30858 :           if (POINTER_TYPE_P (TREE_TYPE (loop->v)))
     362                 :            :             iter_type = long_long_unsigned_type_node;
     363                 :      28890 :           else if (TYPE_UNSIGNED (TREE_TYPE (loop->v))
     364                 :       5344 :                    && TYPE_PRECISION (TREE_TYPE (loop->v))
     365                 :      34234 :                       >= TYPE_PRECISION (iter_type))
     366                 :            :             {
     367                 :       2302 :               tree n;
     368                 :            : 
     369                 :       2302 :               if (loop->cond_code == LT_EXPR)
     370                 :        426 :                 n = fold_build2_loc (loc, PLUS_EXPR, TREE_TYPE (loop->v),
     371                 :            :                                      loop->n2, loop->step);
     372                 :            :               else
     373                 :       1876 :                 n = loop->n1;
     374                 :       2302 :               if (TREE_CODE (n) != INTEGER_CST
     375                 :       2302 :                   || tree_int_cst_lt (TYPE_MAX_VALUE (iter_type), n))
     376                 :       2214 :                 iter_type = long_long_unsigned_type_node;
     377                 :            :             }
     378                 :      26588 :           else if (TYPE_PRECISION (TREE_TYPE (loop->v))
     379                 :      26588 :                    > TYPE_PRECISION (iter_type))
     380                 :            :             {
     381                 :          0 :               tree n1, n2;
     382                 :            : 
     383                 :          0 :               if (loop->cond_code == LT_EXPR)
     384                 :            :                 {
     385                 :          0 :                   n1 = loop->n1;
     386                 :          0 :                   n2 = fold_build2_loc (loc, PLUS_EXPR, TREE_TYPE (loop->v),
     387                 :            :                                         loop->n2, loop->step);
     388                 :            :                 }
     389                 :            :               else
     390                 :            :                 {
     391                 :          0 :                   n1 = fold_build2_loc (loc, MINUS_EXPR, TREE_TYPE (loop->v),
     392                 :            :                                         loop->n2, loop->step);
     393                 :          0 :                   n2 = loop->n1;
     394                 :            :                 }
     395                 :          0 :               if (TREE_CODE (n1) != INTEGER_CST
     396                 :          0 :                   || TREE_CODE (n2) != INTEGER_CST
     397                 :          0 :                   || !tree_int_cst_lt (TYPE_MIN_VALUE (iter_type), n1)
     398                 :          0 :                   || !tree_int_cst_lt (n2, TYPE_MAX_VALUE (iter_type)))
     399                 :          0 :                 iter_type = long_long_unsigned_type_node;
     400                 :            :             }
     401                 :            :         }
     402                 :            : 
     403                 :     156148 :       if (i >= fd->collapse)
     404                 :       1193 :         continue;
     405                 :            : 
     406                 :     154955 :       if (collapse_count && *collapse_count == NULL)
     407                 :            :         {
     408                 :      27447 :           t = fold_binary (loop->cond_code, boolean_type_node,
     409                 :            :                            fold_convert (TREE_TYPE (loop->v), loop->n1),
     410                 :            :                            fold_convert (TREE_TYPE (loop->v), loop->n2));
     411                 :      27447 :           if (t && integer_zerop (t))
     412                 :       3342 :             count = build_zero_cst (long_long_unsigned_type_node);
     413                 :      24105 :           else if ((i == 0 || count != NULL_TREE)
     414                 :      15507 :                    && TREE_CODE (TREE_TYPE (loop->v)) == INTEGER_TYPE
     415                 :      14339 :                    && TREE_CONSTANT (loop->n1)
     416                 :      10652 :                    && TREE_CONSTANT (loop->n2)
     417                 :      34475 :                    && TREE_CODE (loop->step) == INTEGER_CST)
     418                 :            :             {
     419                 :      10368 :               tree itype = TREE_TYPE (loop->v);
     420                 :            : 
     421                 :      10368 :               if (POINTER_TYPE_P (itype))
     422                 :            :                 itype = signed_type_for (itype);
     423                 :      12656 :               t = build_int_cst (itype, (loop->cond_code == LT_EXPR ? -1 : 1));
     424                 :      10368 :               t = fold_build2_loc (loc, PLUS_EXPR, itype,
     425                 :            :                                    fold_convert_loc (loc, itype, loop->step),
     426                 :            :                                    t);
     427                 :      10368 :               t = fold_build2_loc (loc, PLUS_EXPR, itype, t,
     428                 :            :                                    fold_convert_loc (loc, itype, loop->n2));
     429                 :      10368 :               t = fold_build2_loc (loc, MINUS_EXPR, itype, t,
     430                 :            :                                    fold_convert_loc (loc, itype, loop->n1));
     431                 :      10368 :               if (TYPE_UNSIGNED (itype) && loop->cond_code == GT_EXPR)
     432                 :            :                 {
     433                 :       2260 :                   tree step = fold_convert_loc (loc, itype, loop->step);
     434                 :       2260 :                   t = fold_build2_loc (loc, TRUNC_DIV_EXPR, itype,
     435                 :            :                                        fold_build1_loc (loc, NEGATE_EXPR,
     436                 :            :                                                         itype, t),
     437                 :            :                                        fold_build1_loc (loc, NEGATE_EXPR,
     438                 :            :                                                         itype, step));
     439                 :            :                 }
     440                 :            :               else
     441                 :       8108 :                 t = fold_build2_loc (loc, TRUNC_DIV_EXPR, itype, t,
     442                 :            :                                      fold_convert_loc (loc, itype,
     443                 :            :                                                        loop->step));
     444                 :      10368 :               t = fold_convert_loc (loc, long_long_unsigned_type_node, t);
     445                 :      10368 :               if (count != NULL_TREE)
     446                 :       4607 :                 count = fold_build2_loc (loc, MULT_EXPR,
     447                 :            :                                          long_long_unsigned_type_node,
     448                 :            :                                          count, t);
     449                 :            :               else
     450                 :            :                 count = t;
     451                 :      10368 :               if (TREE_CODE (count) != INTEGER_CST)
     452                 :          0 :                 count = NULL_TREE;
     453                 :            :             }
     454                 :      13737 :           else if (count && !integer_zerop (count))
     455                 :            :             count = NULL_TREE;
     456                 :            :         }
     457                 :            :     }
     458                 :            : 
     459                 :      99230 :   if (count
     460                 :      99230 :       && !simd
     461                 :       3501 :       && (fd->sched_kind != OMP_CLAUSE_SCHEDULE_STATIC
     462                 :       2858 :           || fd->have_ordered))
     463                 :            :     {
     464                 :        716 :       if (!tree_int_cst_lt (count, TYPE_MAX_VALUE (long_integer_type_node)))
     465                 :          0 :         iter_type = long_long_unsigned_type_node;
     466                 :            :       else
     467                 :        716 :         iter_type = long_integer_type_node;
     468                 :            :     }
     469                 :      98514 :   else if (collapse_iter && *collapse_iter != NULL)
     470                 :      19853 :     iter_type = TREE_TYPE (*collapse_iter);
     471                 :      99230 :   fd->iter_type = iter_type;
     472                 :      99230 :   if (collapse_iter && *collapse_iter == NULL)
     473                 :       9650 :     *collapse_iter = create_tmp_var (iter_type, ".iter");
     474                 :      99230 :   if (collapse_count && *collapse_count == NULL)
     475                 :            :     {
     476                 :       9650 :       if (count)
     477                 :       4511 :         *collapse_count = fold_convert_loc (loc, iter_type, count);
     478                 :            :       else
     479                 :       5139 :         *collapse_count = create_tmp_var (iter_type, ".count");
     480                 :            :     }
     481                 :            : 
     482                 :      99230 :   if (fd->collapse > 1 || fd->tiling || (fd->ordered && loops))
     483                 :            :     {
     484                 :      29498 :       fd->loop.v = *collapse_iter;
     485                 :      29498 :       fd->loop.n1 = build_int_cst (TREE_TYPE (fd->loop.v), 0);
     486                 :      29498 :       fd->loop.n2 = *collapse_count;
     487                 :      29498 :       fd->loop.step = build_int_cst (TREE_TYPE (fd->loop.v), 1);
     488                 :      29498 :       fd->loop.cond_code = LT_EXPR;
     489                 :            :     }
     490                 :      69732 :   else if (loops)
     491                 :      26723 :     loops[0] = fd->loop;
     492                 :      99230 : }
     493                 :            : 
     494                 :            : /* Build a call to GOMP_barrier.  */
     495                 :            : 
     496                 :            : gimple *
     497                 :       3493 : omp_build_barrier (tree lhs)
     498                 :            : {
     499                 :       3493 :   tree fndecl = builtin_decl_explicit (lhs ? BUILT_IN_GOMP_BARRIER_CANCEL
     500                 :            :                                            : BUILT_IN_GOMP_BARRIER);
     501                 :       3493 :   gcall *g = gimple_build_call (fndecl, 0);
     502                 :       3493 :   if (lhs)
     503                 :         38 :     gimple_call_set_lhs (g, lhs);
     504                 :       3493 :   return g;
     505                 :            : }
     506                 :            : 
     507                 :            : /* Return maximum possible vectorization factor for the target.  */
     508                 :            : 
     509                 :            : poly_uint64
     510                 :      20862 : omp_max_vf (void)
     511                 :            : {
     512                 :      20862 :   if (!optimize
     513                 :      20400 :       || optimize_debug
     514                 :      20400 :       || !flag_tree_loop_optimize
     515                 :      20399 :       || (!flag_tree_loop_vectorize
     516                 :       2565 :           && global_options_set.x_flag_tree_loop_vectorize))
     517                 :        467 :     return 1;
     518                 :            : 
     519                 :      20395 :   auto_vector_modes modes;
     520                 :      20395 :   targetm.vectorize.autovectorize_vector_modes (&modes, true);
     521                 :      20395 :   if (!modes.is_empty ())
     522                 :            :     {
     523                 :      64188 :       poly_uint64 vf = 0;
     524                 :      64188 :       for (unsigned int i = 0; i < modes.length (); ++i)
     525                 :            :         /* The returned modes use the smallest element size (and thus
     526                 :            :            the largest nunits) for the vectorization approach that they
     527                 :            :            represent.  */
     528                 :      87592 :         vf = ordered_max (vf, GET_MODE_NUNITS (modes[i]));
     529                 :      20392 :       return vf;
     530                 :            :     }
     531                 :            : 
     532                 :          3 :   machine_mode vqimode = targetm.vectorize.preferred_simd_mode (QImode);
     533                 :          3 :   if (GET_MODE_CLASS (vqimode) == MODE_VECTOR_INT)
     534                 :          0 :     return GET_MODE_NUNITS (vqimode);
     535                 :            : 
     536                 :          3 :   return 1;
     537                 :            : }
     538                 :            : 
     539                 :            : /* Return maximum SIMT width if offloading may target SIMT hardware.  */
     540                 :            : 
     541                 :            : int
     542                 :       3025 : omp_max_simt_vf (void)
     543                 :            : {
     544                 :       3025 :   if (!optimize)
     545                 :            :     return 0;
     546                 :            :   if (ENABLE_OFFLOADING)
     547                 :            :     for (const char *c = getenv ("OFFLOAD_TARGET_NAMES"); c;)
     548                 :            :       {
     549                 :            :         if (!strncmp (c, "nvptx", strlen ("nvptx")))
     550                 :            :           return 32;
     551                 :            :         else if ((c = strchr (c, ':')))
     552                 :            :           c++;
     553                 :            :       }
     554                 :            :   return 0;
     555                 :            : }
     556                 :            : 
     557                 :            : /* Store the construct selectors as tree codes from last to first,
     558                 :            :    return their number.  */
     559                 :            : 
     560                 :            : int
     561                 :        194 : omp_constructor_traits_to_codes (tree ctx, enum tree_code *constructs)
     562                 :            : {
     563                 :        194 :   int nconstructs = list_length (ctx);
     564                 :        194 :   int i = nconstructs - 1;
     565                 :        580 :   for (tree t2 = ctx; t2; t2 = TREE_CHAIN (t2), i--)
     566                 :            :     {
     567                 :        386 :       const char *sel = IDENTIFIER_POINTER (TREE_PURPOSE (t2));
     568                 :        386 :       if (!strcmp (sel, "target"))
     569                 :          2 :         constructs[i] = OMP_TARGET;
     570                 :        384 :       else if (!strcmp (sel, "teams"))
     571                 :         44 :         constructs[i] = OMP_TEAMS;
     572                 :        340 :       else if (!strcmp (sel, "parallel"))
     573                 :        120 :         constructs[i] = OMP_PARALLEL;
     574                 :        220 :       else if (!strcmp (sel, "for") || !strcmp (sel, "do"))
     575                 :        172 :         constructs[i] = OMP_FOR;
     576                 :         48 :       else if (!strcmp (sel, "simd"))
     577                 :         48 :         constructs[i] = OMP_SIMD;
     578                 :            :       else
     579                 :          0 :         gcc_unreachable ();
     580                 :            :     }
     581                 :        194 :   gcc_assert (i == -1);
     582                 :        194 :   return nconstructs;
     583                 :            : }
     584                 :            : 
     585                 :            : /* Return true if PROP is possibly present in one of the offloading target's
     586                 :            :    OpenMP contexts.  The format of PROPS string is always offloading target's
     587                 :            :    name terminated by '\0', followed by properties for that offloading
     588                 :            :    target separated by '\0' and terminated by another '\0'.  The strings
     589                 :            :    are created from omp-device-properties installed files of all configured
     590                 :            :    offloading targets.  */
     591                 :            : 
     592                 :            : static bool
     593                 :          0 : omp_offload_device_kind_arch_isa (const char *props, const char *prop)
     594                 :            : {
     595                 :          0 :   const char *names = getenv ("OFFLOAD_TARGET_NAMES");
     596                 :          0 :   if (names == NULL || *names == '\0')
     597                 :            :     return false;
     598                 :          0 :   while (*props != '\0')
     599                 :            :     {
     600                 :          0 :       size_t name_len = strlen (props);
     601                 :          0 :       bool matches = false;
     602                 :          0 :       for (const char *c = names; c; )
     603                 :            :         {
     604                 :          0 :           if (strncmp (props, c, name_len) == 0
     605                 :          0 :               && (c[name_len] == '\0'
     606                 :          0 :                   || c[name_len] == ':'
     607                 :          0 :                   || c[name_len] == '='))
     608                 :            :             {
     609                 :            :               matches = true;
     610                 :            :               break;
     611                 :            :             }
     612                 :          0 :           else if ((c = strchr (c, ':')))
     613                 :          0 :             c++;
     614                 :            :         }
     615                 :          0 :       props = props + name_len + 1;
     616                 :          0 :       while (*props != '\0')
     617                 :            :         {
     618                 :          0 :           if (matches && strcmp (props, prop) == 0)
     619                 :            :             return true;
     620                 :          0 :           props = strchr (props, '\0') + 1;
     621                 :            :         }
     622                 :          0 :       props++;
     623                 :            :     }
     624                 :            :   return false;
     625                 :            : }
     626                 :            : 
     627                 :            : /* Return true if the current code location is or might be offloaded.
     628                 :            :    Return true in declare target functions, or when nested in a target
     629                 :            :    region or when unsure, return false otherwise.  */
     630                 :            : 
     631                 :            : static bool
     632                 :        557 : omp_maybe_offloaded (void)
     633                 :            : {
     634                 :        557 :   if (!hsa_gen_requested_p ())
     635                 :            :     {
     636                 :        557 :       if (!ENABLE_OFFLOADING)
     637                 :        467 :         return false;
     638                 :            :       const char *names = getenv ("OFFLOAD_TARGET_NAMES");
     639                 :            :       if (names == NULL || *names == '\0')
     640                 :            :         return false;
     641                 :            :     }
     642                 :            :   if (symtab->state == PARSING)
     643                 :            :     /* Maybe.  */
     644                 :            :     return true;
     645                 :            :   if (current_function_decl
     646                 :            :       && lookup_attribute ("omp declare target",
     647                 :            :                            DECL_ATTRIBUTES (current_function_decl)))
     648                 :            :     return true;
     649                 :            :   if (cfun && (cfun->curr_properties & PROP_gimple_any) == 0)
     650                 :            :     {
     651                 :            :       enum tree_code construct = OMP_TARGET;
     652                 :            :       if (omp_construct_selector_matches (&construct, 1, NULL))
     653                 :            :         return true;
     654                 :            :     }
     655                 :            :   return false;
     656                 :            : }
     657                 :            : 
     658                 :            : /* Return a name from PROP, a property in selectors accepting
     659                 :            :    name lists.  */
     660                 :            : 
     661                 :            : static const char *
     662                 :       1304 : omp_context_name_list_prop (tree prop)
     663                 :            : {
     664                 :       1304 :   if (TREE_PURPOSE (prop))
     665                 :        970 :     return IDENTIFIER_POINTER (TREE_PURPOSE (prop));
     666                 :            :   else
     667                 :            :     {
     668                 :        334 :       const char *ret = TREE_STRING_POINTER (TREE_VALUE (prop));
     669                 :        334 :       if ((size_t) TREE_STRING_LENGTH (TREE_VALUE (prop)) == strlen (ret) + 1)
     670                 :            :         return ret;
     671                 :          2 :       return NULL;
     672                 :            :     }
     673                 :            : }
     674                 :            : 
     675                 :            : /* Return 1 if context selector matches the current OpenMP context, 0
     676                 :            :    if it does not and -1 if it is unknown and need to be determined later.
     677                 :            :    Some properties can be checked right away during parsing (this routine),
     678                 :            :    others need to wait until the whole TU is parsed, others need to wait until
     679                 :            :    IPA, others until vectorization.  */
     680                 :            : 
     681                 :            : int
     682                 :       1483 : omp_context_selector_matches (tree ctx)
     683                 :            : {
     684                 :       1483 :   int ret = 1;
     685                 :       2653 :   for (tree t1 = ctx; t1; t1 = TREE_CHAIN (t1))
     686                 :            :     {
     687                 :       1646 :       char set = IDENTIFIER_POINTER (TREE_PURPOSE (t1))[0];
     688                 :       1646 :       if (set == 'c')
     689                 :            :         {
     690                 :            :           /* For now, ignore the construct set.  While something can be
     691                 :            :              determined already during parsing, we don't know until end of TU
     692                 :            :              whether additional constructs aren't added through declare variant
     693                 :            :              unless "omp declare variant variant" attribute exists already
     694                 :            :              (so in most of the cases), and we'd need to maintain set of
     695                 :            :              surrounding OpenMP constructs, which is better handled during
     696                 :            :              gimplification.  */
     697                 :        335 :           if (symtab->state == PARSING
     698                 :        137 :               || (cfun->curr_properties & PROP_gimple_any) != 0)
     699                 :            :             {
     700                 :        198 :               ret = -1;
     701                 :        302 :               continue;
     702                 :            :             }
     703                 :            : 
     704                 :        137 :           enum tree_code constructs[5];
     705                 :        137 :           int nconstructs
     706                 :        137 :             = omp_constructor_traits_to_codes (TREE_VALUE (t1), constructs);
     707                 :        137 :           int r = omp_construct_selector_matches (constructs, nconstructs,
     708                 :            :                                                   NULL);
     709                 :        137 :           if (r == 0)
     710                 :         33 :             return 0;
     711                 :        104 :           if (r == -1)
     712                 :         20 :             ret = -1;
     713                 :        104 :           continue;
     714                 :            :         }
     715                 :       2329 :       for (tree t2 = TREE_VALUE (t1); t2; t2 = TREE_CHAIN (t2))
     716                 :            :         {
     717                 :       1461 :           const char *sel = IDENTIFIER_POINTER (TREE_PURPOSE (t2));
     718                 :       1461 :           switch (*sel)
     719                 :            :             {
     720                 :        264 :             case 'v':
     721                 :        264 :               if (set == 'i' && !strcmp (sel, "vendor"))
     722                 :        531 :                 for (tree t3 = TREE_VALUE (t2); t3; t3 = TREE_CHAIN (t3))
     723                 :            :                   {
     724                 :        404 :                     const char *prop = omp_context_name_list_prop (t3);
     725                 :        404 :                     if (prop == NULL)
     726                 :            :                       return 0;
     727                 :        125 :                     if ((!strcmp (prop, " score") && TREE_PURPOSE (t3))
     728                 :        403 :                         || !strcmp (prop, "gnu"))
     729                 :        267 :                       continue;
     730                 :            :                     return 0;
     731                 :            :                   }
     732                 :            :               break;
     733                 :         30 :             case 'e':
     734                 :         30 :               if (set == 'i' && !strcmp (sel, "extension"))
     735                 :            :                 /* We don't support any extensions right now.  */
     736                 :            :                 return 0;
     737                 :            :               break;
     738                 :        226 :             case 'a':
     739                 :        226 :               if (set == 'i' && !strcmp (sel, "atomic_default_mem_order"))
     740                 :            :                 {
     741                 :        109 :                   enum omp_memory_order omo
     742                 :            :                     = ((enum omp_memory_order)
     743                 :        109 :                        (omp_requires_mask
     744                 :        109 :                         & OMP_REQUIRES_ATOMIC_DEFAULT_MEM_ORDER));
     745                 :        109 :                   if (omo == OMP_MEMORY_ORDER_UNSPECIFIED)
     746                 :            :                     {
     747                 :            :                       /* We don't know yet, until end of TU.  */
     748                 :         57 :                       if (symtab->state == PARSING)
     749                 :            :                         {
     750                 :            :                           ret = -1;
     751                 :            :                           break;
     752                 :            :                         }
     753                 :            :                       else
     754                 :            :                         omo = OMP_MEMORY_ORDER_RELAXED;
     755                 :            :                     }
     756                 :         52 :                   tree t3 = TREE_VALUE (t2);
     757                 :         52 :                   const char *prop = IDENTIFIER_POINTER (TREE_PURPOSE (t3));
     758                 :         52 :                   if (!strcmp (prop, " score"))
     759                 :            :                     {
     760                 :         50 :                       t3 = TREE_CHAIN (t3);
     761                 :         50 :                       prop = IDENTIFIER_POINTER (TREE_PURPOSE (t3));
     762                 :            :                     }
     763                 :         52 :                   if (!strcmp (prop, "relaxed")
     764                 :          1 :                       && omo != OMP_MEMORY_ORDER_RELAXED)
     765                 :            :                     return 0;
     766                 :         51 :                   else if (!strcmp (prop, "seq_cst")
     767                 :         51 :                            && omo != OMP_MEMORY_ORDER_SEQ_CST)
     768                 :            :                     return 0;
     769                 :         51 :                   else if (!strcmp (prop, "acq_rel")
     770                 :          0 :                            && omo != OMP_MEMORY_ORDER_ACQ_REL)
     771                 :            :                     return 0;
     772                 :            :                 }
     773                 :        168 :               if (set == 'd' && !strcmp (sel, "arch"))
     774                 :        194 :                 for (tree t3 = TREE_VALUE (t2); t3; t3 = TREE_CHAIN (t3))
     775                 :            :                   {
     776                 :        122 :                     const char *arch = omp_context_name_list_prop (t3);
     777                 :        122 :                     if (arch == NULL)
     778                 :            :                       return 0;
     779                 :        122 :                     int r = 0;
     780                 :        122 :                     if (targetm.omp.device_kind_arch_isa != NULL)
     781                 :        122 :                       r = targetm.omp.device_kind_arch_isa (omp_device_arch,
     782                 :            :                                                             arch);
     783                 :        122 :                     if (r == 0 || (r == -1 && symtab->state != PARSING))
     784                 :            :                       {
     785                 :            :                         /* If we are or might be in a target region or
     786                 :            :                            declare target function, need to take into account
     787                 :            :                            also offloading values.  */
     788                 :         45 :                         if (!omp_maybe_offloaded ())
     789                 :         45 :                           return 0;
     790                 :            :                         if (strcmp (arch, "hsa") == 0
     791                 :            :                             && hsa_gen_requested_p ())
     792                 :            :                           {
     793                 :            :                             ret = -1;
     794                 :            :                             continue;
     795                 :            :                           }
     796                 :            :                         if (ENABLE_OFFLOADING)
     797                 :            :                           {
     798                 :            :                             const char *arches = omp_offload_device_arch;
     799                 :            :                             if (omp_offload_device_kind_arch_isa (arches,
     800                 :            :                                                                   arch))
     801                 :            :                               {
     802                 :            :                                 ret = -1;
     803                 :            :                                 continue;
     804                 :            :                               }
     805                 :            :                           }
     806                 :            :                         return 0;
     807                 :            :                       }
     808                 :         77 :                     else if (r == -1)
     809                 :            :                       ret = -1;
     810                 :            :                     /* If arch matches on the host, it still might not match
     811                 :            :                        in the offloading region.  */
     812                 :         77 :                     else if (omp_maybe_offloaded ())
     813                 :            :                       ret = -1;
     814                 :            :                   }
     815                 :            :               break;
     816                 :         35 :             case 'u':
     817                 :         35 :               if (set == 'i' && !strcmp (sel, "unified_address"))
     818                 :            :                 {
     819                 :         20 :                   if ((omp_requires_mask & OMP_REQUIRES_UNIFIED_ADDRESS) == 0)
     820                 :            :                     {
     821                 :         20 :                       if (symtab->state == PARSING)
     822                 :            :                         ret = -1;
     823                 :            :                       else
     824                 :            :                         return 0;
     825                 :            :                     }
     826                 :            :                   break;
     827                 :            :                 }
     828                 :         15 :               if (set == 'i' && !strcmp (sel, "unified_shared_memory"))
     829                 :            :                 {
     830                 :         15 :                   if ((omp_requires_mask
     831                 :         15 :                        & OMP_REQUIRES_UNIFIED_SHARED_MEMORY) == 0)
     832                 :            :                     {
     833                 :         15 :                       if (symtab->state == PARSING)
     834                 :            :                         ret = -1;
     835                 :            :                       else
     836                 :            :                         return 0;
     837                 :            :                     }
     838                 :            :                   break;
     839                 :            :                 }
     840                 :            :               break;
     841                 :         10 :             case 'd':
     842                 :         10 :               if (set == 'i' && !strcmp (sel, "dynamic_allocators"))
     843                 :            :                 {
     844                 :         10 :                   if ((omp_requires_mask
     845                 :         10 :                        & OMP_REQUIRES_DYNAMIC_ALLOCATORS) == 0)
     846                 :            :                     {
     847                 :         10 :                       if (symtab->state == PARSING)
     848                 :            :                         ret = -1;
     849                 :            :                       else
     850                 :            :                         return 0;
     851                 :            :                     }
     852                 :            :                   break;
     853                 :            :                 }
     854                 :            :               break;
     855                 :         10 :             case 'r':
     856                 :         10 :               if (set == 'i' && !strcmp (sel, "reverse_offload"))
     857                 :            :                 {
     858                 :         10 :                   if ((omp_requires_mask & OMP_REQUIRES_REVERSE_OFFLOAD) == 0)
     859                 :            :                     {
     860                 :         10 :                       if (symtab->state == PARSING)
     861                 :            :                         ret = -1;
     862                 :            :                       else
     863                 :            :                         return 0;
     864                 :            :                     }
     865                 :            :                   break;
     866                 :            :                 }
     867                 :            :               break;
     868                 :        203 :             case 'k':
     869                 :        203 :               if (set == 'd' && !strcmp (sel, "kind"))
     870                 :        400 :                 for (tree t3 = TREE_VALUE (t2); t3; t3 = TREE_CHAIN (t3))
     871                 :            :                   {
     872                 :        258 :                     const char *prop = omp_context_name_list_prop (t3);
     873                 :        258 :                     if (prop == NULL)
     874                 :            :                       return 0;
     875                 :        257 :                     if (!strcmp (prop, "any"))
     876                 :         87 :                       continue;
     877                 :        170 :                     if (!strcmp (prop, "host"))
     878                 :            :                       {
     879                 :         25 :                         if (omp_maybe_offloaded ())
     880                 :            :                           ret = -1;
     881                 :         25 :                         continue;
     882                 :            :                       }
     883                 :        145 :                     if (!strcmp (prop, "nohost"))
     884                 :            :                       {
     885                 :            :                         if (omp_maybe_offloaded ())
     886                 :            :                           ret = -1;
     887                 :            :                         else
     888                 :            :                           return 0;
     889                 :            :                         continue;
     890                 :            :                       }
     891                 :        130 :                     int r = 0;
     892                 :        130 :                     if (targetm.omp.device_kind_arch_isa != NULL)
     893                 :        130 :                       r = targetm.omp.device_kind_arch_isa (omp_device_kind,
     894                 :            :                                                             prop);
     895                 :            :                     else
     896                 :          0 :                       r = strcmp (prop, "cpu") == 0;
     897                 :        130 :                     if (r == 0 || (r == -1 && symtab->state != PARSING))
     898                 :            :                       {
     899                 :            :                         /* If we are or might be in a target region or
     900                 :            :                            declare target function, need to take into account
     901                 :            :                            also offloading values.  */
     902                 :            :                         if (!omp_maybe_offloaded ())
     903                 :            :                           return 0;
     904                 :            :                         if (strcmp (prop, "gpu") == 0
     905                 :            :                             && hsa_gen_requested_p ())
     906                 :            :                           {
     907                 :            :                             ret = -1;
     908                 :            :                             continue;
     909                 :            :                           }
     910                 :            :                         if (ENABLE_OFFLOADING)
     911                 :            :                           {
     912                 :            :                             const char *kinds = omp_offload_device_kind;
     913                 :            :                             if (omp_offload_device_kind_arch_isa (kinds, prop))
     914                 :            :                               {
     915                 :            :                                 ret = -1;
     916                 :            :                                 continue;
     917                 :            :                               }
     918                 :            :                           }
     919                 :            :                         return 0;
     920                 :            :                       }
     921                 :         85 :                     else if (r == -1)
     922                 :            :                       ret = -1;
     923                 :            :                     /* If kind matches on the host, it still might not match
     924                 :            :                        in the offloading region.  */
     925                 :         85 :                     else if (omp_maybe_offloaded ())
     926                 :            :                       ret = -1;
     927                 :            :                   }
     928                 :            :               break;
     929                 :        220 :             case 'i':
     930                 :        220 :               if (set == 'd' && !strcmp (sel, "isa"))
     931                 :        610 :                 for (tree t3 = TREE_VALUE (t2); t3; t3 = TREE_CHAIN (t3))
     932                 :            :                   {
     933                 :        410 :                     const char *isa = omp_context_name_list_prop (t3);
     934                 :        410 :                     if (isa == NULL)
     935                 :            :                       return 0;
     936                 :        410 :                     int r = 0;
     937                 :        410 :                     if (targetm.omp.device_kind_arch_isa != NULL)
     938                 :        410 :                       r = targetm.omp.device_kind_arch_isa (omp_device_isa,
     939                 :            :                                                             isa);
     940                 :        410 :                     if (r == 0 || (r == -1 && symtab->state != PARSING))
     941                 :            :                       {
     942                 :            :                         /* If isa is valid on the target, but not in the
     943                 :            :                            current function and current function has
     944                 :            :                            #pragma omp declare simd on it, some simd clones
     945                 :            :                            might have the isa added later on.  */
     946                 :         25 :                         if (r == -1
     947                 :         15 :                             && targetm.simd_clone.compute_vecsize_and_simdlen)
     948                 :            :                           {
     949                 :         15 :                             tree attrs
     950                 :         15 :                               = DECL_ATTRIBUTES (current_function_decl);
     951                 :         15 :                             if (lookup_attribute ("omp declare simd", attrs))
     952                 :            :                               {
     953                 :          5 :                                 ret = -1;
     954                 :          5 :                                 continue;
     955                 :            :                               }
     956                 :            :                           }
     957                 :            :                         /* If we are or might be in a target region or
     958                 :            :                            declare target function, need to take into account
     959                 :            :                            also offloading values.  */
     960                 :         20 :                         if (!omp_maybe_offloaded ())
     961                 :         20 :                           return 0;
     962                 :            :                         if (ENABLE_OFFLOADING)
     963                 :            :                           {
     964                 :            :                             const char *isas = omp_offload_device_isa;
     965                 :            :                             if (omp_offload_device_kind_arch_isa (isas, isa))
     966                 :            :                               {
     967                 :            :                                 ret = -1;
     968                 :            :                                 continue;
     969                 :            :                               }
     970                 :            :                           }
     971                 :            :                         return 0;
     972                 :            :                       }
     973                 :        385 :                     else if (r == -1)
     974                 :            :                       ret = -1;
     975                 :            :                     /* If isa matches on the host, it still might not match
     976                 :            :                        in the offloading region.  */
     977                 :        305 :                     else if (omp_maybe_offloaded ())
     978                 :            :                       ret = -1;
     979                 :            :                   }
     980                 :            :               break;
     981                 :        453 :             case 'c':
     982                 :        453 :               if (set == 'u' && !strcmp (sel, "condition"))
     983                 :        677 :                 for (tree t3 = TREE_VALUE (t2); t3; t3 = TREE_CHAIN (t3))
     984                 :        661 :                   if (TREE_PURPOSE (t3) == NULL_TREE)
     985                 :            :                     {
     986                 :        449 :                       if (integer_zerop (TREE_VALUE (t3)))
     987                 :            :                         return 0;
     988                 :        300 :                       if (integer_nonzerop (TREE_VALUE (t3)))
     989                 :            :                         break;
     990                 :            :                       ret = -1;
     991                 :            :                     }
     992                 :            :               break;
     993                 :            :             default:
     994                 :            :               break;
     995                 :            :             }
     996                 :            :         }
     997                 :            :     }
     998                 :            :   return ret;
     999                 :            : }
    1000                 :            : 
    1001                 :            : /* Compare construct={simd} CLAUSES1 with CLAUSES2, return 0/-1/1/2 as
    1002                 :            :    in omp_context_selector_set_compare.  */
    1003                 :            : 
    1004                 :            : static int
    1005                 :         10 : omp_construct_simd_compare (tree clauses1, tree clauses2)
    1006                 :            : {
    1007                 :         10 :   if (clauses1 == NULL_TREE)
    1008                 :          0 :     return clauses2 == NULL_TREE ? 0 : -1;
    1009                 :         10 :   if (clauses2 == NULL_TREE)
    1010                 :            :     return 1;
    1011                 :            : 
    1012                 :         30 :   int r = 0;
    1013                 :            :   struct declare_variant_simd_data {
    1014                 :            :     bool inbranch, notinbranch;
    1015                 :            :     tree simdlen;
    1016                 :            :     auto_vec<tree,16> data_sharing;
    1017                 :            :     auto_vec<tree,16> aligned;
    1018                 :         20 :     declare_variant_simd_data ()
    1019                 :         20 :       : inbranch(false), notinbranch(false), simdlen(NULL_TREE) {}
    1020                 :         60 :   } data[2];
    1021                 :            :   unsigned int i;
    1022                 :         30 :   for (i = 0; i < 2; i++)
    1023                 :        108 :     for (tree c = i ? clauses2 : clauses1; c; c = OMP_CLAUSE_CHAIN (c))
    1024                 :            :       {
    1025                 :         68 :         vec<tree> *v;
    1026                 :         68 :         switch (OMP_CLAUSE_CODE (c))
    1027                 :            :           {
    1028                 :          8 :           case OMP_CLAUSE_INBRANCH:
    1029                 :          8 :             data[i].inbranch = true;
    1030                 :          8 :             continue;
    1031                 :         12 :           case OMP_CLAUSE_NOTINBRANCH:
    1032                 :         12 :             data[i].notinbranch = true;
    1033                 :         12 :             continue;
    1034                 :         20 :           case OMP_CLAUSE_SIMDLEN:
    1035                 :         20 :             data[i].simdlen = OMP_CLAUSE_SIMDLEN_EXPR (c);
    1036                 :         20 :             continue;
    1037                 :         20 :           case OMP_CLAUSE_UNIFORM:
    1038                 :         20 :           case OMP_CLAUSE_LINEAR:
    1039                 :         20 :             v = &data[i].data_sharing;
    1040                 :         20 :             break;
    1041                 :          8 :           case OMP_CLAUSE_ALIGNED:
    1042                 :          8 :             v = &data[i].aligned;
    1043                 :          8 :             break;
    1044                 :          0 :           default:
    1045                 :          0 :             gcc_unreachable ();
    1046                 :            :           }
    1047                 :         28 :         unsigned HOST_WIDE_INT argno = tree_to_uhwi (OMP_CLAUSE_DECL (c));
    1048                 :         56 :         if (argno >= v->length ())
    1049                 :         28 :           v->safe_grow_cleared (argno + 1);
    1050                 :         28 :         (*v)[argno] = c;
    1051                 :            :       }
    1052                 :            :   /* Here, r is used as a bitmask, 2 is set if CLAUSES1 has something
    1053                 :            :      CLAUSES2 doesn't, 1 is set if CLAUSES2 has something CLAUSES1
    1054                 :            :      doesn't.  Thus, r == 3 implies return value 2, r == 1 implies
    1055                 :            :      -1, r == 2 implies 1 and r == 0 implies 0.  */
    1056                 :         10 :   if (data[0].inbranch != data[1].inbranch)
    1057                 :          0 :     r |= data[0].inbranch ? 2 : 1;
    1058                 :         10 :   if (data[0].notinbranch != data[1].notinbranch)
    1059                 :          0 :     r |= data[0].notinbranch ? 2 : 1;
    1060                 :         10 :   if (!simple_cst_equal (data[0].simdlen, data[1].simdlen))
    1061                 :            :     {
    1062                 :          0 :       if (data[0].simdlen && data[1].simdlen)
    1063                 :            :         return 2;
    1064                 :          0 :       r |= data[0].simdlen ? 2 : 1;
    1065                 :            :     }
    1066                 :         20 :   if (data[0].data_sharing.length () < data[1].data_sharing.length ()
    1067                 :         30 :       || data[0].aligned.length () < data[1].aligned.length ())
    1068                 :          0 :     r |= 1;
    1069                 :            :   tree c1, c2;
    1070                 :         33 :   FOR_EACH_VEC_ELT (data[0].data_sharing, i, c1)
    1071                 :            :     {
    1072                 :         26 :       c2 = (i < data[1].data_sharing.length ()
    1073                 :         26 :             ? data[1].data_sharing[i] : NULL_TREE);
    1074                 :         26 :       if ((c1 == NULL_TREE) != (c2 == NULL_TREE))
    1075                 :            :         {
    1076                 :          0 :           r |= c1 != NULL_TREE ? 2 : 1;
    1077                 :          0 :           continue;
    1078                 :            :         }
    1079                 :         26 :       if (c1 == NULL_TREE)
    1080                 :         16 :         continue;
    1081                 :         10 :       if (OMP_CLAUSE_CODE (c1) != OMP_CLAUSE_CODE (c2))
    1082                 :            :         return 2;
    1083                 :          8 :       if (OMP_CLAUSE_CODE (c1) != OMP_CLAUSE_LINEAR)
    1084                 :          4 :         continue;
    1085                 :          4 :       if (OMP_CLAUSE_LINEAR_VARIABLE_STRIDE (c1)
    1086                 :          4 :           != OMP_CLAUSE_LINEAR_VARIABLE_STRIDE (c2))
    1087                 :            :         return 2;
    1088                 :          4 :       if (OMP_CLAUSE_LINEAR_KIND (c1) != OMP_CLAUSE_LINEAR_KIND (c2))
    1089                 :            :         return 2;
    1090                 :          4 :       if (!simple_cst_equal (OMP_CLAUSE_LINEAR_STEP (c1),
    1091                 :          4 :                              OMP_CLAUSE_LINEAR_STEP (c2)))
    1092                 :            :         return 2;
    1093                 :            :     }
    1094                 :         14 :   FOR_EACH_VEC_ELT (data[0].aligned, i, c1)
    1095                 :            :     {
    1096                 :         18 :       c2 = i < data[1].aligned.length () ? data[1].aligned[i] : NULL_TREE;
    1097                 :          9 :       if ((c1 == NULL_TREE) != (c2 == NULL_TREE))
    1098                 :            :         {
    1099                 :          0 :           r |= c1 != NULL_TREE ? 2 : 1;
    1100                 :          0 :           continue;
    1101                 :            :         }
    1102                 :          9 :       if (c1 == NULL_TREE)
    1103                 :          6 :         continue;
    1104                 :          3 :       if (!simple_cst_equal (OMP_CLAUSE_ALIGNED_ALIGNMENT (c1),
    1105                 :          3 :                              OMP_CLAUSE_ALIGNED_ALIGNMENT (c2)))
    1106                 :            :         return 2;
    1107                 :            :     }
    1108                 :          5 :   switch (r)
    1109                 :            :     {
    1110                 :            :     case 0: return 0;
    1111                 :            :     case 1: return -1;
    1112                 :            :     case 2: return 1;
    1113                 :            :     case 3: return 2;
    1114                 :          0 :     default: gcc_unreachable ();
    1115                 :            :     }
    1116                 :            : }
    1117                 :            : 
    1118                 :            : /* Compare properties of selectors SEL from SET other than construct.
    1119                 :            :    Return 0/-1/1/2 as in omp_context_selector_set_compare.
    1120                 :            :    Unlike set names or selector names, properties can have duplicates.  */
    1121                 :            : 
    1122                 :            : static int
    1123                 :         60 : omp_context_selector_props_compare (const char *set, const char *sel,
    1124                 :            :                                     tree ctx1, tree ctx2)
    1125                 :            : {
    1126                 :         60 :   int ret = 0;
    1127                 :        141 :   for (int pass = 0; pass < 2; pass++)
    1128                 :        295 :     for (tree t1 = pass ? ctx2 : ctx1; t1; t1 = TREE_CHAIN (t1))
    1129                 :            :       {
    1130                 :        115 :         tree t2;
    1131                 :        310 :         for (t2 = pass ? ctx1 : ctx2; t2; t2 = TREE_CHAIN (t2))
    1132                 :        180 :           if (TREE_PURPOSE (t1) == TREE_PURPOSE (t2))
    1133                 :            :             {
    1134                 :         90 :               if (TREE_PURPOSE (t1) == NULL_TREE)
    1135                 :            :                 {
    1136                 :         30 :                   if (set[0] == 'u' && strcmp (sel, "condition") == 0)
    1137                 :            :                     {
    1138                 :         30 :                       if (integer_zerop (TREE_VALUE (t1))
    1139                 :         15 :                           != integer_zerop (TREE_VALUE (t2)))
    1140                 :            :                         return 2;
    1141                 :            :                       break;
    1142                 :            :                     }
    1143                 :         15 :                   if (simple_cst_equal (TREE_VALUE (t1), TREE_VALUE (t2)))
    1144                 :            :                     break;
    1145                 :            :                 }
    1146                 :         60 :               else if (strcmp (IDENTIFIER_POINTER (TREE_PURPOSE (t1)),
    1147                 :            :                                " score") == 0)
    1148                 :            :                 {
    1149                 :         15 :                   if (!simple_cst_equal (TREE_VALUE (t1), TREE_VALUE (t2)))
    1150                 :            :                     return 2;
    1151                 :            :                   break;
    1152                 :            :                 }
    1153                 :            :               else
    1154                 :            :                 break;
    1155                 :            :             }
    1156                 :         90 :           else if (TREE_PURPOSE (t1)
    1157                 :         60 :                    && TREE_PURPOSE (t2) == NULL_TREE
    1158                 :        125 :                    && TREE_CODE (TREE_VALUE (t2)) == STRING_CST)
    1159                 :            :             {
    1160                 :         30 :               const char *p1 = omp_context_name_list_prop (t1);
    1161                 :         30 :               const char *p2 = omp_context_name_list_prop (t2);
    1162                 :         30 :               if (p2
    1163                 :         30 :                   && strcmp (p1, p2) == 0
    1164                 :         10 :                   && strcmp (p1, " score"))
    1165                 :            :                 break;
    1166                 :            :             }
    1167                 :         60 :           else if (TREE_PURPOSE (t1) == NULL_TREE
    1168                 :         30 :                    && TREE_PURPOSE (t2)
    1169                 :         90 :                    && TREE_CODE (TREE_VALUE (t1)) == STRING_CST)
    1170                 :            :             {
    1171                 :         25 :               const char *p1 = omp_context_name_list_prop (t1);
    1172                 :         25 :               const char *p2 = omp_context_name_list_prop (t2);
    1173                 :         25 :               if (p1
    1174                 :         25 :                   && strcmp (p1, p2) == 0
    1175                 :         10 :                   && strcmp (p1, " score"))
    1176                 :            :                 break;
    1177                 :            :             }
    1178                 :        100 :         if (t2 == NULL_TREE)
    1179                 :            :           {
    1180                 :         15 :             int r = pass ? -1 : 1;
    1181                 :         15 :             if (ret && ret != r)
    1182                 :            :               return 2;
    1183                 :         15 :             else if (pass)
    1184                 :          9 :               return r;
    1185                 :            :             else
    1186                 :            :               {
    1187                 :            :                 ret = r;
    1188                 :            :                 break;
    1189                 :            :               }
    1190                 :            :           }
    1191                 :            :       }
    1192                 :            :   return ret;
    1193                 :            : }
    1194                 :            : 
    1195                 :            : /* Compare single context selector sets CTX1 and CTX2 with SET name.
    1196                 :            :    Return 0 if CTX1 is equal to CTX2,
    1197                 :            :    -1 if CTX1 is a strict subset of CTX2,
    1198                 :            :    1 if CTX2 is a strict subset of CTX1, or
    1199                 :            :    2 if neither context is a subset of another one.  */
    1200                 :            : 
    1201                 :            : int
    1202                 :        145 : omp_context_selector_set_compare (const char *set, tree ctx1, tree ctx2)
    1203                 :            : {
    1204                 :        145 :   bool swapped = false;
    1205                 :        145 :   int ret = 0;
    1206                 :        145 :   int len1 = list_length (ctx1);
    1207                 :        145 :   int len2 = list_length (ctx2);
    1208                 :        145 :   int cnt = 0;
    1209                 :        145 :   if (len1 < len2)
    1210                 :            :     {
    1211                 :         35 :       swapped = true;
    1212                 :         35 :       std::swap (ctx1, ctx2);
    1213                 :         35 :       std::swap (len1, len2);
    1214                 :            :     }
    1215                 :        145 :   if (set[0] == 'c')
    1216                 :            :     {
    1217                 :         80 :       tree t1;
    1218                 :         80 :       tree t2 = ctx2;
    1219                 :         80 :       tree simd = get_identifier ("simd");
    1220                 :            :       /* Handle construct set specially.  In this case the order
    1221                 :            :          of the selector matters too.  */
    1222                 :        194 :       for (t1 = ctx1; t1; t1 = TREE_CHAIN (t1))
    1223                 :        189 :         if (TREE_PURPOSE (t1) == TREE_PURPOSE (t2))
    1224                 :            :           {
    1225                 :        130 :             int r = 0;
    1226                 :        130 :             if (TREE_PURPOSE (t1) == simd)
    1227                 :         10 :               r = omp_construct_simd_compare (TREE_VALUE (t1),
    1228                 :         10 :                                               TREE_VALUE (t2));
    1229                 :        130 :             if (r == 2 || (ret && r && (ret < 0) != (r < 0)))
    1230                 :          5 :               return 2;
    1231                 :        125 :             if (ret == 0)
    1232                 :         70 :               ret = r;
    1233                 :        125 :             t2 = TREE_CHAIN (t2);
    1234                 :        125 :             if (t2 == NULL_TREE)
    1235                 :            :               {
    1236                 :         70 :                 t1 = TREE_CHAIN (t1);
    1237                 :         70 :                 break;
    1238                 :            :               }
    1239                 :            :           }
    1240                 :         59 :         else if (ret < 0)
    1241                 :            :           return 2;
    1242                 :            :         else
    1243                 :            :           ret = 1;
    1244                 :         75 :       if (t2 != NULL_TREE)
    1245                 :            :         return 2;
    1246                 :         70 :       if (t1 != NULL_TREE)
    1247                 :            :         {
    1248                 :         20 :           if (ret < 0)
    1249                 :            :             return 2;
    1250                 :            :           ret = 1;
    1251                 :            :         }
    1252                 :         50 :       if (ret == 0)
    1253                 :            :         return 0;
    1254                 :         55 :       return swapped ? -ret : ret;
    1255                 :            :     }
    1256                 :        135 :   for (tree t1 = ctx1; t1; t1 = TREE_CHAIN (t1))
    1257                 :            :     {
    1258                 :            :       tree t2;
    1259                 :        120 :       for (t2 = ctx2; t2; t2 = TREE_CHAIN (t2))
    1260                 :         95 :         if (TREE_PURPOSE (t1) == TREE_PURPOSE (t2))
    1261                 :            :           {
    1262                 :         60 :             const char *sel = IDENTIFIER_POINTER (TREE_PURPOSE (t1));
    1263                 :         60 :             int r = omp_context_selector_props_compare (set, sel,
    1264                 :         60 :                                                         TREE_VALUE (t1),
    1265                 :         60 :                                                         TREE_VALUE (t2));
    1266                 :         60 :             if (r == 2 || (ret && r && (ret < 0) != (r < 0)))
    1267                 :            :               return 2;
    1268                 :         45 :             if (ret == 0)
    1269                 :         40 :               ret = r;
    1270                 :         45 :             cnt++;
    1271                 :         45 :             break;
    1272                 :            :           }
    1273                 :         70 :       if (t2 == NULL_TREE)
    1274                 :            :         {
    1275                 :         25 :           if (ret == -1)
    1276                 :            :             return 2;
    1277                 :            :           ret = 1;
    1278                 :            :         }
    1279                 :            :     }
    1280                 :         50 :   if (cnt < len2)
    1281                 :            :     return 2;
    1282                 :         40 :   if (ret == 0)
    1283                 :            :     return 0;
    1284                 :         25 :   return swapped ? -ret : ret;
    1285                 :            : }
    1286                 :            : 
    1287                 :            : /* Compare whole context selector specification CTX1 and CTX2.
    1288                 :            :    Return 0 if CTX1 is equal to CTX2,
    1289                 :            :    -1 if CTX1 is a strict subset of CTX2,
    1290                 :            :    1 if CTX2 is a strict subset of CTX1, or
    1291                 :            :    2 if neither context is a subset of another one.  */
    1292                 :            : 
    1293                 :            : static int
    1294                 :        217 : omp_context_selector_compare (tree ctx1, tree ctx2)
    1295                 :            : {
    1296                 :        217 :   bool swapped = false;
    1297                 :        217 :   int ret = 0;
    1298                 :        217 :   int len1 = list_length (ctx1);
    1299                 :        217 :   int len2 = list_length (ctx2);
    1300                 :        217 :   int cnt = 0;
    1301                 :        217 :   if (len1 < len2)
    1302                 :            :     {
    1303                 :         33 :       swapped = true;
    1304                 :         33 :       std::swap (ctx1, ctx2);
    1305                 :         33 :       std::swap (len1, len2);
    1306                 :            :     }
    1307                 :        464 :   for (tree t1 = ctx1; t1; t1 = TREE_CHAIN (t1))
    1308                 :            :     {
    1309                 :            :       tree t2;
    1310                 :        479 :       for (t2 = ctx2; t2; t2 = TREE_CHAIN (t2))
    1311                 :        297 :         if (TREE_PURPOSE (t1) == TREE_PURPOSE (t2))
    1312                 :            :           {
    1313                 :        105 :             const char *set = IDENTIFIER_POINTER (TREE_PURPOSE (t1));
    1314                 :        105 :             int r = omp_context_selector_set_compare (set, TREE_VALUE (t1),
    1315                 :        105 :                                                       TREE_VALUE (t2));
    1316                 :        105 :             if (r == 2 || (ret && r && (ret < 0) != (r < 0)))
    1317                 :            :               return 2;
    1318                 :         80 :             if (ret == 0)
    1319                 :         75 :               ret = r;
    1320                 :         80 :             cnt++;
    1321                 :         80 :             break;
    1322                 :            :           }
    1323                 :        262 :       if (t2 == NULL_TREE)
    1324                 :            :         {
    1325                 :        182 :           if (ret == -1)
    1326                 :            :             return 2;
    1327                 :            :           ret = 1;
    1328                 :            :         }
    1329                 :            :     }
    1330                 :        177 :   if (cnt < len2)
    1331                 :            :     return 2;
    1332                 :         60 :   if (ret == 0)
    1333                 :            :     return 0;
    1334                 :         60 :   return swapped ? -ret : ret;
    1335                 :            : }
    1336                 :            : 
    1337                 :            : /* From context selector CTX, return trait-selector with name SEL in
    1338                 :            :    trait-selector-set with name SET if any, or NULL_TREE if not found.
    1339                 :            :    If SEL is NULL, return the list of trait-selectors in SET.  */
    1340                 :            : 
    1341                 :            : tree
    1342                 :       3138 : omp_get_context_selector (tree ctx, const char *set, const char *sel)
    1343                 :            : {
    1344                 :       3138 :   tree setid = get_identifier (set);
    1345                 :       3138 :   tree selid = sel ? get_identifier (sel) : NULL_TREE;
    1346                 :       6229 :   for (tree t1 = ctx; t1; t1 = TREE_CHAIN (t1))
    1347                 :       3489 :     if (TREE_PURPOSE (t1) == setid)
    1348                 :            :       {
    1349                 :        689 :         if (sel == NULL)
    1350                 :        263 :           return TREE_VALUE (t1);
    1351                 :        976 :         for (tree t2 = TREE_VALUE (t1); t2; t2 = TREE_CHAIN (t2))
    1352                 :        685 :           if (TREE_PURPOSE (t2) == selid)
    1353                 :        135 :             return t2;
    1354                 :            :       }
    1355                 :            :   return NULL_TREE;
    1356                 :            : }
    1357                 :            : 
    1358                 :            : /* Compute *SCORE for context selector CTX.  Return true if the score
    1359                 :            :    would be different depending on whether it is a declare simd clone or
    1360                 :            :    not.  DECLARE_SIMD should be true for the case when it would be
    1361                 :            :    a declare simd clone.  */
    1362                 :            : 
    1363                 :            : static bool
    1364                 :        180 : omp_context_compute_score (tree ctx, widest_int *score, bool declare_simd)
    1365                 :            : {
    1366                 :        180 :   tree construct = omp_get_context_selector (ctx, "construct", NULL);
    1367                 :        180 :   bool has_kind = omp_get_context_selector (ctx, "device", "kind");
    1368                 :        180 :   bool has_arch = omp_get_context_selector (ctx, "device", "arch");
    1369                 :        180 :   bool has_isa = omp_get_context_selector (ctx, "device", "isa");
    1370                 :        180 :   bool ret = false;
    1371                 :        180 :   *score = 1;
    1372                 :        400 :   for (tree t1 = ctx; t1; t1 = TREE_CHAIN (t1))
    1373                 :        220 :     if (TREE_VALUE (t1) != construct)
    1374                 :        330 :       for (tree t2 = TREE_VALUE (t1); t2; t2 = TREE_CHAIN (t2))
    1375                 :        165 :         if (tree t3 = TREE_VALUE (t2))
    1376                 :        165 :           if (TREE_PURPOSE (t3)
    1377                 :        145 :               && strcmp (IDENTIFIER_POINTER (TREE_PURPOSE (t3)), " score") == 0
    1378                 :        270 :               && TREE_CODE (TREE_VALUE (t3)) == INTEGER_CST)
    1379                 :        105 :             *score += wi::to_widest (TREE_VALUE (t3));
    1380                 :        180 :   if (construct || has_kind || has_arch || has_isa)
    1381                 :            :     {
    1382                 :        105 :       int scores[12];
    1383                 :        105 :       enum tree_code constructs[5];
    1384                 :        105 :       int nconstructs = 0;
    1385                 :        105 :       if (construct)
    1386                 :         55 :         nconstructs = omp_constructor_traits_to_codes (construct, constructs);
    1387                 :        105 :       if (omp_construct_selector_matches (constructs, nconstructs, scores)
    1388                 :            :           == 2)
    1389                 :         30 :         ret = true;
    1390                 :        105 :       int b = declare_simd ? nconstructs + 1 : 0;
    1391                 :        105 :       if (scores[b + nconstructs] + 4U < score->get_precision ())
    1392                 :            :         {
    1393                 :        214 :           for (int n = 0; n < nconstructs; ++n)
    1394                 :            :             {
    1395                 :        109 :               if (scores[b + n] < 0)
    1396                 :            :                 {
    1397                 :          0 :                   *score = -1;
    1398                 :          0 :                   return ret;
    1399                 :            :                 }
    1400                 :        109 :               *score += wi::shifted_mask <widest_int> (scores[b + n], 1, false);
    1401                 :            :             }
    1402                 :        105 :           if (has_kind)
    1403                 :         20 :             *score += wi::shifted_mask <widest_int> (scores[b + nconstructs],
    1404                 :         20 :                                                      1, false);
    1405                 :        105 :           if (has_arch)
    1406                 :          5 :             *score += wi::shifted_mask <widest_int> (scores[b + nconstructs] + 1,
    1407                 :          5 :                                                      1, false);
    1408                 :        105 :           if (has_isa)
    1409                 :         25 :             *score += wi::shifted_mask <widest_int> (scores[b + nconstructs] + 2,
    1410                 :         25 :                                                      1, false);
    1411                 :            :         }
    1412                 :            :       else /* FIXME: Implement this.  */
    1413                 :          0 :         gcc_unreachable ();
    1414                 :            :     }
    1415                 :            :   return ret;
    1416                 :            : }
    1417                 :            : 
    1418                 :            : /* Try to resolve declare variant, return the variant decl if it should
    1419                 :            :    be used instead of base, or base otherwise.  */
    1420                 :            : 
    1421                 :            : tree
    1422                 :     125730 : omp_resolve_declare_variant (tree base)
    1423                 :            : {
    1424                 :     125730 :   tree variant1 = NULL_TREE, variant2 = NULL_TREE;
    1425                 :     125730 :   auto_vec <tree, 16> variants;
    1426                 :     125730 :   auto_vec <bool, 16> defer;
    1427                 :     125730 :   bool any_deferred = false;
    1428                 :     126111 :   for (tree attr = DECL_ATTRIBUTES (base); attr; attr = TREE_CHAIN (attr))
    1429                 :            :     {
    1430                 :      56513 :       attr = lookup_attribute ("omp declare variant base", attr);
    1431                 :      56513 :       if (attr == NULL_TREE)
    1432                 :            :         break;
    1433                 :        381 :       if (TREE_CODE (TREE_PURPOSE (TREE_VALUE (attr))) != FUNCTION_DECL)
    1434                 :          0 :         continue;
    1435                 :        381 :       switch (omp_context_selector_matches (TREE_VALUE (TREE_VALUE (attr))))
    1436                 :            :         {
    1437                 :            :         case 0:
    1438                 :            :           /* No match, ignore.  */
    1439                 :            :           break;
    1440                 :         25 :         case -1:
    1441                 :            :           /* Needs to be deferred.  */
    1442                 :         25 :           any_deferred = true;
    1443                 :         25 :           variants.safe_push (attr);
    1444                 :         25 :           defer.safe_push (true);
    1445                 :         25 :           break;
    1446                 :        312 :         default:
    1447                 :        312 :           variants.safe_push (attr);
    1448                 :        312 :           defer.safe_push (false);
    1449                 :        312 :           break;
    1450                 :            :         }
    1451                 :            :     }
    1452                 :     125730 :   if (variants.length () == 0)
    1453                 :            :     return base;
    1454                 :            : 
    1455                 :        169 :   if (any_deferred)
    1456                 :            :     {
    1457                 :         17 :       widest_int max_score1 = 0;
    1458                 :         17 :       widest_int max_score2 = 0;
    1459                 :         17 :       bool first = true;
    1460                 :         17 :       unsigned int i;
    1461                 :         17 :       tree attr1, attr2;
    1462                 :         57 :       FOR_EACH_VEC_ELT (variants, i, attr1)
    1463                 :            :         {
    1464                 :         40 :           widest_int score1;
    1465                 :         40 :           widest_int score2;
    1466                 :         40 :           bool need_two;
    1467                 :         40 :           tree ctx = TREE_VALUE (TREE_VALUE (attr1));
    1468                 :         40 :           need_two = omp_context_compute_score (ctx, &score1, false);
    1469                 :         40 :           if (need_two)
    1470                 :         15 :             omp_context_compute_score (ctx, &score2, true);
    1471                 :            :           else
    1472                 :         25 :             score2 = score1;
    1473                 :         40 :           if (first)
    1474                 :            :             {
    1475                 :         17 :               first = false;
    1476                 :         17 :               max_score1 = score1;
    1477                 :         17 :               max_score2 = score2;
    1478                 :         17 :               if (!defer[i])
    1479                 :            :                 {
    1480                 :          1 :                   variant1 = attr1;
    1481                 :          1 :                   variant2 = attr1;
    1482                 :            :                 }
    1483                 :            :             }
    1484                 :            :           else
    1485                 :            :             {
    1486                 :         46 :               if (max_score1 == score1)
    1487                 :            :                 variant1 = NULL_TREE;
    1488                 :         19 :               else if (score1 > max_score1)
    1489                 :            :                 {
    1490                 :         13 :                   max_score1 = score1;
    1491                 :         13 :                   variant1 = defer[i] ? NULL_TREE : attr1;
    1492                 :            :                 }
    1493                 :         46 :               if (max_score2 == score2)
    1494                 :            :                 variant2 = NULL_TREE;
    1495                 :         23 :               else if (score2 > max_score2)
    1496                 :            :                 {
    1497                 :         13 :                   max_score2 = score2;
    1498                 :         13 :                   variant2 = defer[i] ? NULL_TREE : attr1;
    1499                 :            :                 }
    1500                 :            :             }
    1501                 :            :         }
    1502                 :            : 
    1503                 :            :       /* If there is a clear winner variant with the score which is not
    1504                 :            :          deferred, verify it is not a strict subset of any other context
    1505                 :            :          selector and if it is not, it is the best alternative no matter
    1506                 :            :          whether the others do or don't match.  */
    1507                 :         17 :       if (variant1 && variant1 == variant2)
    1508                 :            :         {
    1509                 :          5 :           tree ctx1 = TREE_VALUE (TREE_VALUE (variant1));
    1510                 :         25 :           FOR_EACH_VEC_ELT (variants, i, attr2)
    1511                 :            :             {
    1512                 :         20 :               if (attr2 == variant1)
    1513                 :          5 :                 continue;
    1514                 :         15 :               tree ctx2 = TREE_VALUE (TREE_VALUE (attr2));
    1515                 :         15 :               int r = omp_context_selector_compare (ctx1, ctx2);
    1516                 :         15 :               if (r == -1)
    1517                 :            :                 {
    1518                 :            :                   /* The winner is a strict subset of ctx2, can't
    1519                 :            :                      decide now.  */
    1520                 :            :                   variant1 = NULL_TREE;
    1521                 :            :                   break;
    1522                 :            :                 }
    1523                 :            :             }
    1524                 :          5 :           if (variant1)
    1525                 :          5 :             return TREE_PURPOSE (TREE_VALUE (variant1));
    1526                 :            :         }
    1527                 :            : 
    1528                 :         12 :       return base;
    1529                 :            :     }
    1530                 :            : 
    1531                 :        152 :   if (variants.length () == 1)
    1532                 :         82 :     return TREE_PURPOSE (TREE_VALUE (variants[0]));
    1533                 :            : 
    1534                 :            :   /* A context selector that is a strict subset of another context selector has a score
    1535                 :            :      of zero.  */
    1536                 :            :   tree attr1, attr2;
    1537                 :            :   unsigned int i, j;
    1538                 :        285 :   FOR_EACH_VEC_ELT (variants, i, attr1)
    1539                 :        215 :     if (attr1)
    1540                 :            :       {
    1541                 :        184 :         tree ctx1 = TREE_VALUE (TREE_VALUE (attr1));
    1542                 :        574 :         FOR_EACH_VEC_ELT_FROM (variants, j, attr2, i + 1)
    1543                 :        204 :           if (attr2)
    1544                 :            :             {
    1545                 :        202 :               tree ctx2 = TREE_VALUE (TREE_VALUE (attr2));
    1546                 :        202 :               int r = omp_context_selector_compare (ctx1, ctx2);
    1547                 :        202 :               if (r == -1)
    1548                 :            :                 {
    1549                 :            :                   /* ctx1 is a strict subset of ctx2, remove
    1550                 :            :                      attr1 from the vector.  */
    1551                 :         29 :                   variants[i] = NULL_TREE;
    1552                 :         29 :                   break;
    1553                 :            :                 }
    1554                 :        173 :               else if (r == 1)
    1555                 :            :                 /* ctx2 is a strict subset of ctx1, remove attr2
    1556                 :            :                    from the vector.  */
    1557                 :         31 :                 variants[j] = NULL_TREE;
    1558                 :            :             }
    1559                 :            :       }
    1560                 :         70 :   widest_int max_score1 = 0;
    1561                 :         70 :   widest_int max_score2 = 0;
    1562                 :         70 :   bool first = true;
    1563                 :        285 :   FOR_EACH_VEC_ELT (variants, i, attr1)
    1564                 :        215 :     if (attr1)
    1565                 :            :       {
    1566                 :        155 :         if (variant1)
    1567                 :            :           {
    1568                 :         85 :             widest_int score1;
    1569                 :         85 :             widest_int score2;
    1570                 :         85 :             bool need_two;
    1571                 :         85 :             tree ctx;
    1572                 :         85 :             if (first)
    1573                 :            :               {
    1574                 :         40 :                 first = false;
    1575                 :         40 :                 ctx = TREE_VALUE (TREE_VALUE (variant1));
    1576                 :         40 :                 need_two = omp_context_compute_score (ctx, &max_score1, false);
    1577                 :         40 :                 if (need_two)
    1578                 :          0 :                   omp_context_compute_score (ctx, &max_score2, true);
    1579                 :            :                 else
    1580                 :         40 :                   max_score2 = max_score1;
    1581                 :            :               }
    1582                 :         85 :             ctx = TREE_VALUE (TREE_VALUE (attr1));
    1583                 :         85 :             need_two = omp_context_compute_score (ctx, &score1, false);
    1584                 :         85 :             if (need_two)
    1585                 :          0 :               omp_context_compute_score (ctx, &score2, true);
    1586                 :            :             else
    1587                 :         85 :               score2 = score1;
    1588                 :         85 :             if (score1 > max_score1)
    1589                 :            :               {
    1590                 :         26 :                 max_score1 = score1;
    1591                 :         26 :                 variant1 = attr1;
    1592                 :            :               }
    1593                 :         85 :             if (score2 > max_score2)
    1594                 :            :               {
    1595                 :         26 :                 max_score2 = score2;
    1596                 :         26 :                 variant2 = attr1;
    1597                 :            :               }
    1598                 :            :           }
    1599                 :            :         else
    1600                 :            :           {
    1601                 :            :             variant1 = attr1;
    1602                 :            :             variant2 = attr1;
    1603                 :            :           }
    1604                 :            :       }
    1605                 :            :   /* If there is a disagreement on which variant has the highest score
    1606                 :            :      depending on whether it will be in a declare simd clone or not,
    1607                 :            :      punt for now and defer until after IPA where we will know that.  */
    1608                 :         70 :   return ((variant1 && variant1 == variant2)
    1609                 :        140 :           ? TREE_PURPOSE (TREE_VALUE (variant1)) : base);
    1610                 :            : }
    1611                 :            : 
    1612                 :            : 
    1613                 :            : /* Encode an oacc launch argument.  This matches the GOMP_LAUNCH_PACK
    1614                 :            :    macro on gomp-constants.h.  We do not check for overflow.  */
    1615                 :            : 
    1616                 :            : tree
    1617                 :       8383 : oacc_launch_pack (unsigned code, tree device, unsigned op)
    1618                 :            : {
    1619                 :       8383 :   tree res;
    1620                 :            : 
    1621                 :       8383 :   res = build_int_cst (unsigned_type_node, GOMP_LAUNCH_PACK (code, 0, op));
    1622                 :       8383 :   if (device)
    1623                 :            :     {
    1624                 :          0 :       device = fold_build2 (LSHIFT_EXPR, unsigned_type_node,
    1625                 :            :                             device, build_int_cst (unsigned_type_node,
    1626                 :            :                                                    GOMP_LAUNCH_DEVICE_SHIFT));
    1627                 :          0 :       res = fold_build2 (BIT_IOR_EXPR, unsigned_type_node, res, device);
    1628                 :            :     }
    1629                 :       8383 :   return res;
    1630                 :            : }
    1631                 :            : 
    1632                 :            : /* FIXME: What is the following comment for? */
    1633                 :            : /* Look for compute grid dimension clauses and convert to an attribute
    1634                 :            :    attached to FN.  This permits the target-side code to (a) massage
    1635                 :            :    the dimensions, (b) emit that data and (c) optimize.  Non-constant
    1636                 :            :    dimensions are pushed onto ARGS.
    1637                 :            : 
    1638                 :            :    The attribute value is a TREE_LIST.  A set of dimensions is
    1639                 :            :    represented as a list of INTEGER_CST.  Those that are runtime
    1640                 :            :    exprs are represented as an INTEGER_CST of zero.
    1641                 :            : 
    1642                 :            :    TODO: Normally the attribute will just contain a single such list.  If
    1643                 :            :    however it contains a list of lists, this will represent the use of
    1644                 :            :    device_type.  Each member of the outer list is an assoc list of
    1645                 :            :    dimensions, keyed by the device type.  The first entry will be the
    1646                 :            :    default.  Well, that's the plan.  */
    1647                 :            : 
    1648                 :            : /* Replace any existing oacc fn attribute with updated dimensions.  */
    1649                 :            : 
    1650                 :            : /* Variant working on a list of attributes.  */
    1651                 :            : 
    1652                 :            : tree
    1653                 :      16902 : oacc_replace_fn_attrib_attr (tree attribs, tree dims)
    1654                 :            : {
    1655                 :      16902 :   tree ident = get_identifier (OACC_FN_ATTRIB);
    1656                 :            : 
    1657                 :            :   /* If we happen to be present as the first attrib, drop it.  */
    1658                 :      32969 :   if (attribs && TREE_PURPOSE (attribs) == ident)
    1659                 :       8122 :     attribs = TREE_CHAIN (attribs);
    1660                 :      16902 :   return tree_cons (ident, dims, attribs);
    1661                 :            : }
    1662                 :            : 
    1663                 :            : /* Variant working on a function decl.  */
    1664                 :            : 
    1665                 :            : void
    1666                 :      16723 : oacc_replace_fn_attrib (tree fn, tree dims)
    1667                 :            : {
    1668                 :      16723 :   DECL_ATTRIBUTES (fn)
    1669                 :      16723 :     = oacc_replace_fn_attrib_attr (DECL_ATTRIBUTES (fn), dims);
    1670                 :      16723 : }
    1671                 :            : 
    1672                 :            : /* Scan CLAUSES for launch dimensions and attach them to the oacc
    1673                 :            :    function attribute.  Push any that are non-constant onto the ARGS
    1674                 :            :    list, along with an appropriate GOMP_LAUNCH_DIM tag.  */
    1675                 :            : 
    1676                 :            : void
    1677                 :       8464 : oacc_set_fn_attrib (tree fn, tree clauses, vec<tree> *args)
    1678                 :            : {
    1679                 :            :   /* Must match GOMP_DIM ordering.  */
    1680                 :       8464 :   static const omp_clause_code ids[]
    1681                 :            :     = { OMP_CLAUSE_NUM_GANGS, OMP_CLAUSE_NUM_WORKERS,
    1682                 :            :         OMP_CLAUSE_VECTOR_LENGTH };
    1683                 :       8464 :   unsigned ix;
    1684                 :       8464 :   tree dims[GOMP_DIM_MAX];
    1685                 :            : 
    1686                 :       8464 :   tree attr = NULL_TREE;
    1687                 :       8464 :   unsigned non_const = 0;
    1688                 :            : 
    1689                 :      33856 :   for (ix = GOMP_DIM_MAX; ix--;)
    1690                 :            :     {
    1691                 :      25392 :       tree clause = omp_find_clause (clauses, ids[ix]);
    1692                 :      25392 :       tree dim = NULL_TREE;
    1693                 :            : 
    1694                 :      25392 :       if (clause)
    1695                 :       3037 :         dim = OMP_CLAUSE_EXPR (clause, ids[ix]);
    1696                 :      25392 :       dims[ix] = dim;
    1697                 :      25392 :       if (dim && TREE_CODE (dim) != INTEGER_CST)
    1698                 :            :         {
    1699                 :        180 :           dim = integer_zero_node;
    1700                 :        180 :           non_const |= GOMP_DIM_MASK (ix);
    1701                 :            :         }
    1702                 :      25392 :       attr = tree_cons (NULL_TREE, dim, attr);
    1703                 :            :     }
    1704                 :            : 
    1705                 :       8464 :   oacc_replace_fn_attrib (fn, attr);
    1706                 :            : 
    1707                 :       8464 :   if (non_const)
    1708                 :            :     {
    1709                 :            :       /* Push a dynamic argument set.  */
    1710                 :        114 :       args->safe_push (oacc_launch_pack (GOMP_LAUNCH_DIM,
    1711                 :            :                                          NULL_TREE, non_const));
    1712                 :        456 :       for (unsigned ix = 0; ix != GOMP_DIM_MAX; ix++)
    1713                 :        342 :         if (non_const & GOMP_DIM_MASK (ix))
    1714                 :        180 :           args->safe_push (dims[ix]);
    1715                 :            :     }
    1716                 :       8464 : }
    1717                 :            : 
    1718                 :            : /* Verify OpenACC routine clauses.
    1719                 :            : 
    1720                 :            :    Returns 0 if FNDECL should be marked with an OpenACC 'routine' directive, 1
    1721                 :            :    if it has already been marked in compatible way, and -1 if incompatible.
    1722                 :            :    Upon returning, the chain of clauses will contain exactly one clause
    1723                 :            :    specifying the level of parallelism.  */
    1724                 :            : 
    1725                 :            : int
    1726                 :       1108 : oacc_verify_routine_clauses (tree fndecl, tree *clauses, location_t loc,
    1727                 :            :                              const char *routine_str)
    1728                 :            : {
    1729                 :       1108 :   tree c_level = NULL_TREE;
    1730                 :       1108 :   tree c_p = NULL_TREE;
    1731                 :       3819 :   for (tree c = *clauses; c; c_p = c, c = OMP_CLAUSE_CHAIN (c))
    1732                 :       2711 :     switch (OMP_CLAUSE_CODE (c))
    1733                 :            :       {
    1734                 :       2711 :       case OMP_CLAUSE_GANG:
    1735                 :       2711 :       case OMP_CLAUSE_WORKER:
    1736                 :       2711 :       case OMP_CLAUSE_VECTOR:
    1737                 :       2711 :       case OMP_CLAUSE_SEQ:
    1738                 :       2711 :         if (c_level == NULL_TREE)
    1739                 :            :           c_level = c;
    1740                 :       1870 :         else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_CODE (c_level))
    1741                 :            :           {
    1742                 :            :             /* This has already been diagnosed in the front ends.  */
    1743                 :            :             /* Drop the duplicate clause.  */
    1744                 :        440 :             gcc_checking_assert (c_p != NULL_TREE);
    1745                 :        440 :             OMP_CLAUSE_CHAIN (c_p) = OMP_CLAUSE_CHAIN (c);
    1746                 :        440 :             c = c_p;
    1747                 :            :           }
    1748                 :            :         else
    1749                 :            :           {
    1750                 :       1430 :             error_at (OMP_CLAUSE_LOCATION (c),
    1751                 :            :                       "%qs specifies a conflicting level of parallelism",
    1752                 :       1430 :                       omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
    1753                 :       1430 :             inform (OMP_CLAUSE_LOCATION (c_level),
    1754                 :            :                     "... to the previous %qs clause here",
    1755                 :       1430 :                     omp_clause_code_name[OMP_CLAUSE_CODE (c_level)]);
    1756                 :            :             /* Drop the conflicting clause.  */
    1757                 :       1430 :             gcc_checking_assert (c_p != NULL_TREE);
    1758                 :       1430 :             OMP_CLAUSE_CHAIN (c_p) = OMP_CLAUSE_CHAIN (c);
    1759                 :       1430 :             c = c_p;
    1760                 :            :           }
    1761                 :       2711 :         break;
    1762                 :          0 :       default:
    1763                 :          0 :         gcc_unreachable ();
    1764                 :            :       }
    1765                 :       1108 :   if (c_level == NULL_TREE)
    1766                 :            :     {
    1767                 :            :       /* Default to an implicit 'seq' clause.  */
    1768                 :        267 :       c_level = build_omp_clause (loc, OMP_CLAUSE_SEQ);
    1769                 :        267 :       OMP_CLAUSE_CHAIN (c_level) = *clauses;
    1770                 :        267 :       *clauses = c_level;
    1771                 :            :     }
    1772                 :            :   /* In *clauses, we now have exactly one clause specifying the level of
    1773                 :            :      parallelism.  */
    1774                 :            : 
    1775                 :       1108 :   tree attr
    1776                 :       1108 :     = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (fndecl));
    1777                 :       1108 :   if (attr != NULL_TREE)
    1778                 :            :     {
    1779                 :            :       /* If a "#pragma acc routine" has already been applied, just verify
    1780                 :            :          this one for compatibility.  */
    1781                 :            :       /* Collect previous directive's clauses.  */
    1782                 :        390 :       tree c_level_p = NULL_TREE;
    1783                 :        780 :       for (tree c = TREE_VALUE (attr); c; c = OMP_CLAUSE_CHAIN (c))
    1784                 :        390 :         switch (OMP_CLAUSE_CODE (c))
    1785                 :            :           {
    1786                 :        390 :           case OMP_CLAUSE_GANG:
    1787                 :        390 :           case OMP_CLAUSE_WORKER:
    1788                 :        390 :           case OMP_CLAUSE_VECTOR:
    1789                 :        390 :           case OMP_CLAUSE_SEQ:
    1790                 :        390 :             gcc_checking_assert (c_level_p == NULL_TREE);
    1791                 :        390 :             c_level_p = c;
    1792                 :        390 :             break;
    1793                 :          0 :           default:
    1794                 :          0 :             gcc_unreachable ();
    1795                 :            :           }
    1796                 :        390 :       gcc_checking_assert (c_level_p != NULL_TREE);
    1797                 :            :       /* ..., and compare to current directive's, which we've already collected
    1798                 :            :          above.  */
    1799                 :        390 :       tree c_diag;
    1800                 :        390 :       tree c_diag_p;
    1801                 :            :       /* Matching level of parallelism?  */
    1802                 :        390 :       if (OMP_CLAUSE_CODE (c_level) != OMP_CLAUSE_CODE (c_level_p))
    1803                 :            :         {
    1804                 :        120 :           c_diag = c_level;
    1805                 :        120 :           c_diag_p = c_level_p;
    1806                 :        120 :           goto incompatible;
    1807                 :            :         }
    1808                 :            :       /* Compatible.  */
    1809                 :            :       return 1;
    1810                 :            : 
    1811                 :        120 :     incompatible:
    1812                 :        120 :       if (c_diag != NULL_TREE)
    1813                 :        120 :         error_at (OMP_CLAUSE_LOCATION (c_diag),
    1814                 :            :                   "incompatible %qs clause when applying"
    1815                 :            :                   " %<%s%> to %qD, which has already been"
    1816                 :            :                   " marked with an OpenACC 'routine' directive",
    1817                 :        120 :                   omp_clause_code_name[OMP_CLAUSE_CODE (c_diag)],
    1818                 :            :                   routine_str, fndecl);
    1819                 :            :       else if (c_diag_p != NULL_TREE)
    1820                 :            :         error_at (loc,
    1821                 :            :                   "missing %qs clause when applying"
    1822                 :            :                   " %<%s%> to %qD, which has already been"
    1823                 :            :                   " marked with an OpenACC 'routine' directive",
    1824                 :            :                   omp_clause_code_name[OMP_CLAUSE_CODE (c_diag_p)],
    1825                 :            :                   routine_str, fndecl);
    1826                 :            :       else
    1827                 :            :         gcc_unreachable ();
    1828                 :        120 :       if (c_diag_p != NULL_TREE)
    1829                 :        120 :         inform (OMP_CLAUSE_LOCATION (c_diag_p),
    1830                 :            :                 "... with %qs clause here",
    1831                 :        120 :                 omp_clause_code_name[OMP_CLAUSE_CODE (c_diag_p)]);
    1832                 :            :       else
    1833                 :            :         {
    1834                 :            :           /* In the front ends, we don't preserve location information for the
    1835                 :            :              OpenACC routine directive itself.  However, that of c_level_p
    1836                 :            :              should be close.  */
    1837                 :            :           location_t loc_routine = OMP_CLAUSE_LOCATION (c_level_p);
    1838                 :            :           inform (loc_routine, "... without %qs clause near to here",
    1839                 :            :                   omp_clause_code_name[OMP_CLAUSE_CODE (c_diag)]);
    1840                 :            :         }
    1841                 :            :       /* Incompatible.  */
    1842                 :        120 :       return -1;
    1843                 :            :     }
    1844                 :            : 
    1845                 :            :   return 0;
    1846                 :            : }
    1847                 :            : 
    1848                 :            : /*  Process the OpenACC 'routine' directive clauses to generate an attribute
    1849                 :            :     for the level of parallelism.  All dimensions have a size of zero
    1850                 :            :     (dynamic).  TREE_PURPOSE is set to indicate whether that dimension
    1851                 :            :     can have a loop partitioned on it.  non-zero indicates
    1852                 :            :     yes, zero indicates no.  By construction once a non-zero has been
    1853                 :            :     reached, further inner dimensions must also be non-zero.  We set
    1854                 :            :     TREE_VALUE to zero for the dimensions that may be partitioned and
    1855                 :            :     1 for the other ones -- if a loop is (erroneously) spawned at
    1856                 :            :     an outer level, we don't want to try and partition it.  */
    1857                 :            : 
    1858                 :            : tree
    1859                 :        887 : oacc_build_routine_dims (tree clauses)
    1860                 :            : {
    1861                 :            :   /* Must match GOMP_DIM ordering.  */
    1862                 :        887 :   static const omp_clause_code ids[]
    1863                 :            :     = {OMP_CLAUSE_GANG, OMP_CLAUSE_WORKER, OMP_CLAUSE_VECTOR, OMP_CLAUSE_SEQ};
    1864                 :        887 :   int ix;
    1865                 :        887 :   int level = -1;
    1866                 :            : 
    1867                 :       1774 :   for (; clauses; clauses = OMP_CLAUSE_CHAIN (clauses))
    1868                 :       1665 :     for (ix = GOMP_DIM_MAX + 1; ix--;)
    1869                 :       1665 :       if (OMP_CLAUSE_CODE (clauses) == ids[ix])
    1870                 :            :         {
    1871                 :            :           level = ix;
    1872                 :            :           break;
    1873                 :            :         }
    1874                 :        887 :   gcc_checking_assert (level >= 0);
    1875                 :            : 
    1876                 :            :   tree dims = NULL_TREE;
    1877                 :            : 
    1878                 :       3548 :   for (ix = GOMP_DIM_MAX; ix--;)
    1879                 :       2661 :     dims = tree_cons (build_int_cst (boolean_type_node, ix >= level),
    1880                 :       2661 :                       build_int_cst (integer_type_node, ix < level), dims);
    1881                 :            : 
    1882                 :        887 :   return dims;
    1883                 :            : }
    1884                 :            : 
    1885                 :            : /* Retrieve the oacc function attrib and return it.  Non-oacc
    1886                 :            :    functions will return NULL.  */
    1887                 :            : 
    1888                 :            : tree
    1889                 :     171675 : oacc_get_fn_attrib (tree fn)
    1890                 :            : {
    1891                 :     171675 :   return lookup_attribute (OACC_FN_ATTRIB, DECL_ATTRIBUTES (fn));
    1892                 :            : }
    1893                 :            : 
    1894                 :            : /* Return true if FN is an OpenMP or OpenACC offloading function.  */
    1895                 :            : 
    1896                 :            : bool
    1897                 :        316 : offloading_function_p (tree fn)
    1898                 :            : {
    1899                 :        316 :   tree attrs = DECL_ATTRIBUTES (fn);
    1900                 :        316 :   return (lookup_attribute ("omp declare target", attrs)
    1901                 :        316 :           || lookup_attribute ("omp target entrypoint", attrs));
    1902                 :            : }
    1903                 :            : 
    1904                 :            : /* Extract an oacc execution dimension from FN.  FN must be an
    1905                 :            :    offloaded function or routine that has already had its execution
    1906                 :            :    dimensions lowered to the target-specific values.  */
    1907                 :            : 
    1908                 :            : int
    1909                 :      16610 : oacc_get_fn_dim_size (tree fn, int axis)
    1910                 :            : {
    1911                 :      16610 :   tree attrs = oacc_get_fn_attrib (fn);
    1912                 :            : 
    1913                 :      16610 :   gcc_assert (axis < GOMP_DIM_MAX);
    1914                 :            : 
    1915                 :      16610 :   tree dims = TREE_VALUE (attrs);
    1916                 :      34154 :   while (axis--)
    1917                 :      17544 :     dims = TREE_CHAIN (dims);
    1918                 :            : 
    1919                 :      16610 :   int size = TREE_INT_CST_LOW (TREE_VALUE (dims));
    1920                 :            : 
    1921                 :      16610 :   return size;
    1922                 :            : }
    1923                 :            : 
    1924                 :            : /* Extract the dimension axis from an IFN_GOACC_DIM_POS or
    1925                 :            :    IFN_GOACC_DIM_SIZE call.  */
    1926                 :            : 
    1927                 :            : int
    1928                 :      16610 : oacc_get_ifn_dim_arg (const gimple *stmt)
    1929                 :            : {
    1930                 :      16610 :   gcc_checking_assert (gimple_call_internal_fn (stmt) == IFN_GOACC_DIM_SIZE
    1931                 :            :                        || gimple_call_internal_fn (stmt) == IFN_GOACC_DIM_POS);
    1932                 :      16610 :   tree arg = gimple_call_arg (stmt, 0);
    1933                 :      16610 :   HOST_WIDE_INT axis = TREE_INT_CST_LOW (arg);
    1934                 :            : 
    1935                 :      16610 :   gcc_checking_assert (axis >= 0 && axis < GOMP_DIM_MAX);
    1936                 :      16610 :   return (int) axis;
    1937                 :            : }

Generated by: LCOV version 1.0

LCOV profile is generated on x86_64 machine using following configure options: configure --disable-bootstrap --enable-coverage=opt --enable-languages=c,c++,fortran,go,jit,lto --enable-host-shared. GCC test suite is run with the built compiler.