openmp: Add support for 2 argument num_teams clause

In OpenMP 5.1, num_teams clause can accept either one expression as before,
but it in that case changed meaning, rather than create <= expression
teams it is now create == expression teams.  Or it accepts two expressions
separated by :, with the meaning that the first is low bound and second upper
bound on how many teams should be created.  The other ways to set number of
teams are upper bounds with lower bound of 1.

The following patch does parsing of this for C/C++.  For host teams, we
actually don't need to do anything further right now, we always create
(pretend to create) exactly the requested number of teams, so we can just
evaluate and throw away the lower bound for now.
For teams nested in target, we don't guarantee that though and further
work will be needed.
In particular, omplower now turns the teams part of:
struct S { S (); S (const S &); ~S (); int s; };
void bar (S &, S &);
int baz ();
_Pragma ("omp declare target to (baz)");

void
foo (void)
{
  S a, b;
  #pragma omp target private (a) map (b)
  {
    #pragma omp teams firstprivate (b) num_teams (baz ())
    {
      bar (a, b);
    }
  }
}
into:
  retval.0 = baz ();
  retval.1 = retval.0;
  {
    unsigned int retval.3;
    struct S * D.2549;
    struct S b;

    retval.3 = (unsigned int) retval.1;
    D.2549 = .omp_data_i->b;
    S::S (&b, D.2549);
    #pragma omp teams num_teams(retval.1) firstprivate(b) shared(a)
    __builtin_GOMP_teams (retval.3, 0);
    {
      bar (&a, &b);
    }
    S::~S (&b);
    #pragma omp return(nowait)
  }
IMHO we want a new API, say GOMP_teams3 which will take 3 arguments
instead of 2 (the lower and upper bounds from num_teams and thread_limit)
and will return a bool whether it should do the teams body or not.
And, we should add right before outermost {} above
while (__builtin_GOMP_teams3 ((unsigned) retval.1, (unsigned) retval.1, 0))
and remove the __builtin_GOMP_teams call.  The current function performs
exit equivalent (at least on NVPTX) which seems bad because that means
the destructors of e.g. private variables on target aren't invoked, and
at the current placement neither destructors of the already constructed
privatized variables in teams.
I'll do this next on the compiler side, but I'm afraid I'll need help
with the nvptx and amdgcn implementations.  E.g. for nvptx, we won't be
able to use %ctaid.x .  I think ideal would be to use a .shared
integer variable for the omp_get_team_num value, but I don't have any
experience with that, are .shared variables zero initialized by default,
or do they have random value at start?  PTX docs say they aren't initializable.

2021-11-11  Jakub Jelinek  <jakub@redhat.com>

gcc/
	* tree.h (OMP_CLAUSE_NUM_TEAMS_EXPR): Rename to ...
	(OMP_CLAUSE_NUM_TEAMS_UPPER_EXPR): ... this.
	(OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR): Define.
	* tree.c (omp_clause_num_ops): Increase num ops for
	OMP_CLAUSE_NUM_TEAMS to 2.
	* tree-pretty-print.c (dump_omp_clause): Print optional lower bound
	for OMP_CLAUSE_NUM_TEAMS.
	* gimplify.c (gimplify_scan_omp_clauses): Gimplify
	OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR if non-NULL.
	(optimize_target_teams): Use OMP_CLAUSE_NUM_TEAMS_UPPER_EXPR instead
	of OMP_CLAUSE_NUM_TEAMS_EXPR.  Handle OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR.
	* omp-low.c (lower_omp_teams): Use OMP_CLAUSE_NUM_TEAMS_UPPER_EXPR
	instead of OMP_CLAUSE_NUM_TEAMS_EXPR.
	* omp-expand.c (expand_teams_call, get_target_arguments): Likewise.
gcc/c/
	* c-parser.c (c_parser_omp_clause_num_teams): Parse optional
	lower-bound and store it into OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR.
	Use OMP_CLAUSE_NUM_TEAMS_UPPER_EXPR instead of
	OMP_CLAUSE_NUM_TEAMS_EXPR.
	(c_parser_omp_target): For OMP_CLAUSE_NUM_TEAMS evaluate before
	combined target teams even lower-bound expression.
gcc/cp/
	* parser.c (cp_parser_omp_clause_num_teams): Parse optional
	lower-bound and store it into OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR.
	Use OMP_CLAUSE_NUM_TEAMS_UPPER_EXPR instead of
	OMP_CLAUSE_NUM_TEAMS_EXPR.
	(cp_parser_omp_target): For OMP_CLAUSE_NUM_TEAMS evaluate before
	combined target teams even lower-bound expression.
	* semantics.c (finish_omp_clauses): Handle
	OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR of OMP_CLAUSE_NUM_TEAMS clause.
	* pt.c (tsubst_omp_clauses): Likewise.
	(tsubst_expr): For OMP_CLAUSE_NUM_TEAMS evaluate before
	combined target teams even lower-bound expression.
gcc/fortran/
	* trans-openmp.c (gfc_trans_omp_clauses): Use
	OMP_CLAUSE_NUM_TEAMS_UPPER_EXPR instead of OMP_CLAUSE_NUM_TEAMS_EXPR.
gcc/testsuite/
	* c-c++-common/gomp/clauses-1.c (bar): Supply lower-bound expression
	to half of the num_teams clauses.
	* c-c++-common/gomp/num-teams-1.c: New test.
	* c-c++-common/gomp/num-teams-2.c: New test.
	* g++.dg/gomp/attrs-1.C (bar): Supply lower-bound expression
	to half of the num_teams clauses.
	* g++.dg/gomp/attrs-2.C (bar): Likewise.
	* g++.dg/gomp/num-teams-1.C: New test.
	* g++.dg/gomp/num-teams-2.C: New test.
libgomp/
	* testsuite/libgomp.c-c++-common/teams-1.c: New test.

diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 80dd61d599ef49a7c58f87bc93ecdcfc27c56fcf..3c9f5877481fd2a9a33381f1787cbb8bd289845b 100644 (file)
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -15175,7 +15175,10 @@ c_parser_omp_clause_orderedkind (c_parser *parser ATTRIBUTE_UNUSED,
 }
 
 /* OpenMP 4.0:
-   num_teams ( expression ) */
+   num_teams ( expression )
+
+   OpenMP 5.1:
+   num_teams ( expression : expression ) */
 
 static tree
 c_parser_omp_clause_num_teams (c_parser *parser, tree list)
@@ -15184,34 +15187,68 @@ c_parser_omp_clause_num_teams (c_parser *parser, tree list)
   matching_parens parens;
   if (parens.require_open (parser))
     {
-      location_t expr_loc = c_parser_peek_token (parser)->location;
+      location_t upper_loc = c_parser_peek_token (parser)->location;
+      location_t lower_loc = UNKNOWN_LOCATION;
       c_expr expr = c_parser_expr_no_commas (parser, NULL);
-      expr = convert_lvalue_to_rvalue (expr_loc, expr, false, true);
-      tree c, t = expr.value;
-      t = c_fully_fold (t, false, NULL);
+      expr = convert_lvalue_to_rvalue (upper_loc, expr, false, true);
+      tree c, upper = expr.value, lower = NULL_TREE;
+      upper = c_fully_fold (upper, false, NULL);
+
+      if (c_parser_next_token_is (parser, CPP_COLON))
+       {
+         c_parser_consume_token (parser);
+         lower_loc = upper_loc;
+         lower = upper;
+         upper_loc = c_parser_peek_token (parser)->location;
+         expr = c_parser_expr_no_commas (parser, NULL);
+         expr = convert_lvalue_to_rvalue (upper_loc, expr, false, true);
+         upper = expr.value;
+         upper = c_fully_fold (upper, false, NULL);
+       }
 
       parens.skip_until_found_close (parser);
 
-      if (!INTEGRAL_TYPE_P (TREE_TYPE (t)))
+      if (!INTEGRAL_TYPE_P (TREE_TYPE (upper))
+         || (lower && !INTEGRAL_TYPE_P (TREE_TYPE (lower))))
        {
          c_parser_error (parser, "expected integer expression");
          return list;
        }
 
       /* Attempt to statically determine when the number isn't positive.  */
-      c = fold_build2_loc (expr_loc, LE_EXPR, boolean_type_node, t,
-                          build_int_cst (TREE_TYPE (t), 0));
-      protected_set_expr_location (c, expr_loc);
+      c = fold_build2_loc (upper_loc, LE_EXPR, boolean_type_node, upper,
+                          build_int_cst (TREE_TYPE (upper), 0));
+      protected_set_expr_location (c, upper_loc);
       if (c == boolean_true_node)
        {
-         warning_at (expr_loc, 0, "%<num_teams%> value must be positive");
-         t = integer_one_node;
+         warning_at (upper_loc, 0, "%<num_teams%> value must be positive");
+         upper = integer_one_node;
+       }
+      if (lower)
+       {
+         c = fold_build2_loc (lower_loc, LE_EXPR, boolean_type_node, lower,
+                              build_int_cst (TREE_TYPE (lower), 0));
+         protected_set_expr_location (c, lower_loc);
+         if (c == boolean_true_node)
+           {
+             warning_at (lower_loc, 0, "%<num_teams%> value must be positive");
+             lower = NULL_TREE;
+           }
+         else if (TREE_CODE (lower) == INTEGER_CST
+                  && TREE_CODE (upper) == INTEGER_CST
+                  && tree_int_cst_lt (upper, lower))
+           {
+             warning_at (lower_loc, 0, "%<num_teams%> lower bound %qE bigger "
+                                       "than upper bound %qE", lower, upper);
+             lower = NULL_TREE;
+           }
        }
 
       check_no_duplicate_clause (list, OMP_CLAUSE_NUM_TEAMS, "num_teams");
 
       c = build_omp_clause (num_teams_loc, OMP_CLAUSE_NUM_TEAMS);
-      OMP_CLAUSE_NUM_TEAMS_EXPR (c) = t;
+      OMP_CLAUSE_NUM_TEAMS_UPPER_EXPR (c) = upper;
+      OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR (c) = lower;
       OMP_CLAUSE_CHAIN (c) = list;
       list = c;
     }
@@ -21016,31 +21053,31 @@ c_parser_omp_target (c_parser *parser, enum pragma_context context, bool *if_p)
          if (ret == NULL_TREE)
            return false;
          if (ccode == OMP_TEAMS)
-           {
-             /* For combined target teams, ensure the num_teams and
-                thread_limit clause expressions are evaluated on the host,
-                before entering the target construct.  */
-             tree c;
-             for (c = cclauses[C_OMP_CLAUSE_SPLIT_TEAMS];
-                  c; c = OMP_CLAUSE_CHAIN (c))
-               if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NUM_TEAMS
-                    || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_THREAD_LIMIT)
-                   && TREE_CODE (OMP_CLAUSE_OPERAND (c, 0)) != INTEGER_CST)
-                 {
-                   tree expr = OMP_CLAUSE_OPERAND (c, 0);
-                   tree tmp = create_tmp_var_raw (TREE_TYPE (expr));
-                   expr = build4 (TARGET_EXPR, TREE_TYPE (expr), tmp,
-                                  expr, NULL_TREE, NULL_TREE);
-                   add_stmt (expr);
-                   OMP_CLAUSE_OPERAND (c, 0) = expr;
-                   tree tc = build_omp_clause (OMP_CLAUSE_LOCATION (c),
-                                               OMP_CLAUSE_FIRSTPRIVATE);
-                   OMP_CLAUSE_DECL (tc) = tmp;
-                   OMP_CLAUSE_CHAIN (tc)
-                     = cclauses[C_OMP_CLAUSE_SPLIT_TARGET];
-                   cclauses[C_OMP_CLAUSE_SPLIT_TARGET] = tc;
-                 }
-           }
+           /* For combined target teams, ensure the num_teams and
+              thread_limit clause expressions are evaluated on the host,
+              before entering the target construct.  */
+           for (tree c = cclauses[C_OMP_CLAUSE_SPLIT_TEAMS];
+                c; c = OMP_CLAUSE_CHAIN (c))
+             if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NUM_TEAMS
+                 || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_THREAD_LIMIT)
+               for (int i = 0;
+                    i <= (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NUM_TEAMS); ++i)
+                 if (OMP_CLAUSE_OPERAND (c, i)
+                     && TREE_CODE (OMP_CLAUSE_OPERAND (c, i)) != INTEGER_CST)
+                   {
+                     tree expr = OMP_CLAUSE_OPERAND (c, i);
+                     tree tmp = create_tmp_var_raw (TREE_TYPE (expr));
+                     expr = build4 (TARGET_EXPR, TREE_TYPE (expr), tmp,
+                                    expr, NULL_TREE, NULL_TREE);
+                     add_stmt (expr);
+                     OMP_CLAUSE_OPERAND (c, i) = expr;
+                     tree tc = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+                                                 OMP_CLAUSE_FIRSTPRIVATE);
+                     OMP_CLAUSE_DECL (tc) = tmp;
+                     OMP_CLAUSE_CHAIN (tc)
+                       = cclauses[C_OMP_CLAUSE_SPLIT_TARGET];
+                     cclauses[C_OMP_CLAUSE_SPLIT_TARGET] = tc;
+                   }
          tree stmt = make_node (OMP_TARGET);
          TREE_TYPE (stmt) = void_type_node;
          OMP_TARGET_CLAUSES (stmt) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET];

diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 32de97b08bdddac03c905e51091aa2fd889c9c11..4c1fed02c7498ff4305a5ef1fbb989e859b6f94c 100644 (file)
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -38275,21 +38275,35 @@ cp_parser_omp_clause_orderedkind (cp_parser * /*parser*/,
 }
 
 /* OpenMP 4.0:
-   num_teams ( expression ) */
+   num_teams ( expression )
+
+   OpenMP 5.1:
+   num_teams ( expression : expression ) */
 
 static tree
 cp_parser_omp_clause_num_teams (cp_parser *parser, tree list,
                                location_t location)
 {
-  tree t, c;
+  tree upper, lower = NULL_TREE, c;
 
   matching_parens parens;
   if (!parens.require_open (parser))
     return list;
 
-  t = cp_parser_assignment_expression (parser);
+  bool saved_colon_corrects_to_scope_p = parser->colon_corrects_to_scope_p;
+  parser->colon_corrects_to_scope_p = false;
+  upper = cp_parser_assignment_expression (parser);
+  parser->colon_corrects_to_scope_p = saved_colon_corrects_to_scope_p;
 
-  if (t == error_mark_node
+  if (upper != error_mark_node
+      && cp_lexer_next_token_is (parser->lexer, CPP_COLON))
+    {
+      lower = upper;
+      cp_lexer_consume_token (parser->lexer);
+      upper = cp_parser_assignment_expression (parser);
+    }
+
+  if (upper == error_mark_node
       || !parens.require_close (parser))
     cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/true,
                                           /*or_comma=*/false,
@@ -38299,7 +38313,8 @@ cp_parser_omp_clause_num_teams (cp_parser *parser, tree list,
                             "num_teams", location);
 
   c = build_omp_clause (location, OMP_CLAUSE_NUM_TEAMS);
-  OMP_CLAUSE_NUM_TEAMS_EXPR (c) = t;
+  OMP_CLAUSE_NUM_TEAMS_UPPER_EXPR (c) = upper;
+  OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR (c) = lower;
   OMP_CLAUSE_CHAIN (c) = list;
 
   return c;
@@ -44104,32 +44119,33 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
          if (ret == NULL_TREE)
            return false;
          if (ccode == OMP_TEAMS && !processing_template_decl)
-           {
-             /* For combined target teams, ensure the num_teams and
-                thread_limit clause expressions are evaluated on the host,
-                before entering the target construct.  */
-             tree c;
-             for (c = cclauses[C_OMP_CLAUSE_SPLIT_TEAMS];
-                  c; c = OMP_CLAUSE_CHAIN (c))
-               if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NUM_TEAMS
-                    || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_THREAD_LIMIT)
-                   && TREE_CODE (OMP_CLAUSE_OPERAND (c, 0)) != INTEGER_CST)
-                 {
-                   tree expr = OMP_CLAUSE_OPERAND (c, 0);
-                   expr = force_target_expr (TREE_TYPE (expr), expr, tf_none);
-                   if (expr == error_mark_node)
-                     continue;
-                   tree tmp = TARGET_EXPR_SLOT (expr);
-                   add_stmt (expr);
-                   OMP_CLAUSE_OPERAND (c, 0) = expr;
-                   tree tc = build_omp_clause (OMP_CLAUSE_LOCATION (c),
-                                               OMP_CLAUSE_FIRSTPRIVATE);
-                   OMP_CLAUSE_DECL (tc) = tmp;
-                   OMP_CLAUSE_CHAIN (tc)
-                     = cclauses[C_OMP_CLAUSE_SPLIT_TARGET];
-                   cclauses[C_OMP_CLAUSE_SPLIT_TARGET] = tc;
-                 }
-           }
+           /* For combined target teams, ensure the num_teams and
+              thread_limit clause expressions are evaluated on the host,
+              before entering the target construct.  */
+           for (tree c = cclauses[C_OMP_CLAUSE_SPLIT_TEAMS];
+                c; c = OMP_CLAUSE_CHAIN (c))
+             if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NUM_TEAMS
+                 || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_THREAD_LIMIT)
+               for (int i = 0;
+                    i <= (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NUM_TEAMS); ++i)
+                 if (OMP_CLAUSE_OPERAND (c, i)
+                     && TREE_CODE (OMP_CLAUSE_OPERAND (c, i)) != INTEGER_CST)
+                   {
+                     tree expr = OMP_CLAUSE_OPERAND (c, i);
+                     expr = force_target_expr (TREE_TYPE (expr), expr,
+                                               tf_none);
+                     if (expr == error_mark_node)
+                       continue;
+                     tree tmp = TARGET_EXPR_SLOT (expr);
+                     add_stmt (expr);
+                     OMP_CLAUSE_OPERAND (c, i) = expr;
+                     tree tc = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+                                                 OMP_CLAUSE_FIRSTPRIVATE);
+                     OMP_CLAUSE_DECL (tc) = tmp;
+                     OMP_CLAUSE_CHAIN (tc)
+                       = cclauses[C_OMP_CLAUSE_SPLIT_TARGET];
+                     cclauses[C_OMP_CLAUSE_SPLIT_TARGET] = tc;
+                   }
          tree stmt = make_node (OMP_TARGET);
          TREE_TYPE (stmt) = void_type_node;
          OMP_TARGET_CLAUSES (stmt) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET];

diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c
index b2916f8aa4b6ea9a3feba18ced4a9f0c374cfd2c..82bf7dc26f6eba133eda7ace35d8d275fb119417 100644 (file)
--- a/gcc/cp/pt.c
+++ b/gcc/cp/pt.c
@@ -17437,6 +17437,13 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort,
            = tsubst_omp_clause_decl (OMP_CLAUSE_DECL (oc), args, complain,
                                      in_decl, iterator_cache);
          break;
+       case OMP_CLAUSE_NUM_TEAMS:
+         if (OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR (oc))
+           OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR (nc)
+             = tsubst_expr (OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR (oc), args,
+                            complain, in_decl,
+                            /*integral_constant_expression_p=*/false);
+         /* FALLTHRU */
        case OMP_CLAUSE_TILE:
        case OMP_CLAUSE_IF:
        case OMP_CLAUSE_NUM_THREADS:
@@ -17445,7 +17452,6 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort,
        case OMP_CLAUSE_FINAL:
        case OMP_CLAUSE_DEVICE:
        case OMP_CLAUSE_DIST_SCHEDULE:
-       case OMP_CLAUSE_NUM_TEAMS:
        case OMP_CLAUSE_THREAD_LIMIT:
        case OMP_CLAUSE_SAFELEN:
        case OMP_CLAUSE_SIMDLEN:
@@ -18948,31 +18954,32 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
        {
          tree teams = cp_walk_tree (&stmt, tsubst_find_omp_teams, NULL, NULL);
          if (teams)
-           {
-             /* For combined target teams, ensure the num_teams and
-                thread_limit clause expressions are evaluated on the host,
-                before entering the target construct.  */
-             tree c;
-             for (c = OMP_TEAMS_CLAUSES (teams);
-                  c; c = OMP_CLAUSE_CHAIN (c))
-               if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NUM_TEAMS
-                    || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_THREAD_LIMIT)
-                   && TREE_CODE (OMP_CLAUSE_OPERAND (c, 0)) != INTEGER_CST)
-                 {
-                   tree expr = OMP_CLAUSE_OPERAND (c, 0);
-                   expr = force_target_expr (TREE_TYPE (expr), expr, tf_none);
-                   if (expr == error_mark_node)
-                     continue;
-                   tmp = TARGET_EXPR_SLOT (expr);
-                   add_stmt (expr);
-                   OMP_CLAUSE_OPERAND (c, 0) = expr;
-                   tree tc = build_omp_clause (OMP_CLAUSE_LOCATION (c),
-                                               OMP_CLAUSE_FIRSTPRIVATE);
-                   OMP_CLAUSE_DECL (tc) = tmp;
-                   OMP_CLAUSE_CHAIN (tc) = OMP_TARGET_CLAUSES (t);
-                   OMP_TARGET_CLAUSES (t) = tc;
-                 }
-           }
+           /* For combined target teams, ensure the num_teams and
+              thread_limit clause expressions are evaluated on the host,
+              before entering the target construct.  */
+           for (tree c = OMP_TEAMS_CLAUSES (teams);
+                c; c = OMP_CLAUSE_CHAIN (c))
+             if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NUM_TEAMS
+                 || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_THREAD_LIMIT)
+               for (int i = 0;
+                    i <= (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NUM_TEAMS); ++i)
+                 if (OMP_CLAUSE_OPERAND (c, i)
+                     && TREE_CODE (OMP_CLAUSE_OPERAND (c, i)) != INTEGER_CST)
+                   {
+                     tree expr = OMP_CLAUSE_OPERAND (c, i);
+                     expr = force_target_expr (TREE_TYPE (expr), expr,
+                                               tf_none);
+                     if (expr == error_mark_node)
+                       continue;
+                     tmp = TARGET_EXPR_SLOT (expr);
+                     add_stmt (expr);
+                     OMP_CLAUSE_OPERAND (c, i) = expr;
+                     tree tc = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+                                                 OMP_CLAUSE_FIRSTPRIVATE);
+                     OMP_CLAUSE_DECL (tc) = tmp;
+                     OMP_CLAUSE_CHAIN (tc) = OMP_TARGET_CLAUSES (t);
+                     OMP_TARGET_CLAUSES (t) = tc;
+                   }
        }
       add_stmt (t);
       break;

diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 2443d0327498bf2fba11d80b51a4762a945dfe4e..60e0982cc480427f3e835b82eeb9ee00401b3706 100644 (file)
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -7228,6 +7228,53 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
                }
              OMP_CLAUSE_OPERAND (c, 0) = t;
            }
+         if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NUM_TEAMS
+             && OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR (c)
+             && !remove)
+           {
+             t = OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR (c);
+             if (t == error_mark_node)
+               remove = true;
+             else if (!type_dependent_expression_p (t)
+                      && !INTEGRAL_TYPE_P (TREE_TYPE (t)))
+               {
+                 error_at (OMP_CLAUSE_LOCATION (c),
+                           "%qs expression must be integral",
+                           omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+                 remove = true;
+               }
+             else
+               {
+                 t = mark_rvalue_use (t);
+                 if (!processing_template_decl)
+                   {
+                     t = maybe_constant_value (t);
+                     if (TREE_CODE (t) == INTEGER_CST
+                         && tree_int_cst_sgn (t) != 1)
+                       {
+                         warning_at (OMP_CLAUSE_LOCATION (c), 0,
+                                     "%qs value must be positive",
+                                     omp_clause_code_name
+                                     [OMP_CLAUSE_CODE (c)]);
+                         t = NULL_TREE;
+                       }
+                     else
+                       t = fold_build_cleanup_point_expr (TREE_TYPE (t), t);
+                     tree upper = OMP_CLAUSE_NUM_TEAMS_UPPER_EXPR (c);
+                     if (t
+                         && TREE_CODE (t) == INTEGER_CST
+                         && TREE_CODE (upper) == INTEGER_CST
+                         && tree_int_cst_lt (upper, t))
+                       {
+                         warning_at (OMP_CLAUSE_LOCATION (c), 0,
+                                     "%<num_teams%> lower bound %qE bigger "
+                                     "than upper bound %qE", t, upper);
+                         t = NULL_TREE;
+                       }
+                   }
+                 OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR (c) = t;
+               }
+           }
          break;
 
        case OMP_CLAUSE_SCHEDULE:

diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index e81c5588c53c1e20a78f165f558b712cc2d295cb..22d66629c07c8b5138a23f340638257b1bf38faf 100644 (file)
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -3938,7 +3938,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
       gfc_add_block_to_block (block, &se.post);
 
       c = build_omp_clause (gfc_get_location (&where), OMP_CLAUSE_NUM_TEAMS);
-      OMP_CLAUSE_NUM_TEAMS_EXPR (c) = num_teams;
+      OMP_CLAUSE_NUM_TEAMS_UPPER_EXPR (c) = num_teams;
       omp_clauses = gfc_trans_add_clause (c, omp_clauses);
     }
 

diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index c2ab96e7e1822db0353bb5bf852a2af435928181..e5877c8309952e5e72a65a237f89b2d6b9ac8354 100644 (file)
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -10273,9 +10273,24 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
            = gimple_boolify (OMP_CLAUSE_OPERAND (c, 0));
          /* Fall through.  */
 
+       case OMP_CLAUSE_NUM_TEAMS:
+         if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NUM_TEAMS
+             && OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR (c)
+             && !is_gimple_min_invariant (OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR (c)))
+           {
+             if (error_operand_p (OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR (c)))
+               {
+                 remove = true;
+                 break;
+               }
+             OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR (c)
+               = get_initialized_tmp_var (OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR (c),
+                                          pre_p, NULL, true);
+           }
+         /* Fall through.  */
+
        case OMP_CLAUSE_SCHEDULE:
        case OMP_CLAUSE_NUM_THREADS:
-       case OMP_CLAUSE_NUM_TEAMS:
        case OMP_CLAUSE_THREAD_LIMIT:
        case OMP_CLAUSE_DIST_SCHEDULE:
        case OMP_CLAUSE_DEVICE:
@@ -13535,7 +13550,8 @@ optimize_target_teams (tree target, gimple_seq *pre_p)
 {
   tree body = OMP_BODY (target);
   tree teams = walk_tree (&body, find_omp_teams, NULL, NULL);
-  tree num_teams = integer_zero_node;
+  tree num_teams_lower = NULL_TREE;
+  tree num_teams_upper = integer_zero_node;
   tree thread_limit = integer_zero_node;
   location_t num_teams_loc = EXPR_LOCATION (target);
   location_t thread_limit_loc = EXPR_LOCATION (target);
@@ -13543,14 +13559,42 @@ optimize_target_teams (tree target, gimple_seq *pre_p)
   struct gimplify_omp_ctx *target_ctx = gimplify_omp_ctxp;
 
   if (teams == NULL_TREE)
-    num_teams = integer_one_node;
+    num_teams_upper = integer_one_node;
   else
     for (c = OMP_TEAMS_CLAUSES (teams); c; c = OMP_CLAUSE_CHAIN (c))
       {
        if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NUM_TEAMS)
          {
-           p = &num_teams;
+           p = &num_teams_upper;
            num_teams_loc = OMP_CLAUSE_LOCATION (c);
+           if (OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR (c))
+             {
+               expr = OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR (c);
+               if (TREE_CODE (expr) == INTEGER_CST)
+                 num_teams_lower = expr;
+               else if (walk_tree (&expr, computable_teams_clause,
+                                   NULL, NULL))
+                 num_teams_lower = integer_minus_one_node;
+               else
+                 {
+                   num_teams_lower = expr;
+                   gimplify_omp_ctxp = gimplify_omp_ctxp->outer_context;
+                   if (gimplify_expr (&num_teams_lower, pre_p, NULL,
+                                      is_gimple_val, fb_rvalue, false)
+                       == GS_ERROR)
+                     {
+                       gimplify_omp_ctxp = target_ctx;
+                       num_teams_lower = integer_minus_one_node;
+                     }
+                   else
+                     {
+                       gimplify_omp_ctxp = target_ctx;
+                       if (!DECL_P (expr) && TREE_CODE (expr) != TARGET_EXPR)
+                         OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR (c)
+                           = num_teams_lower;
+                     }
+                 }
+             }
          }
        else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_THREAD_LIMIT)
          {
@@ -13588,7 +13632,8 @@ optimize_target_teams (tree target, gimple_seq *pre_p)
   OMP_CLAUSE_CHAIN (c) = OMP_TARGET_CLAUSES (target);
   OMP_TARGET_CLAUSES (target) = c;
   c = build_omp_clause (num_teams_loc, OMP_CLAUSE_NUM_TEAMS);
-  OMP_CLAUSE_NUM_TEAMS_EXPR (c) = num_teams;
+  OMP_CLAUSE_NUM_TEAMS_UPPER_EXPR (c) = num_teams_upper;
+  OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR (c) = num_teams_lower;
   OMP_CLAUSE_CHAIN (c) = OMP_TARGET_CLAUSES (target);
   OMP_TARGET_CLAUSES (target) = c;
 }

diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c
index 70957a66da834bf2d86965163001d45bbd71882c..1a5426e35170d981cf07f2cff649985896ddec48 100644 (file)
--- a/gcc/omp-expand.c
+++ b/gcc/omp-expand.c
@@ -938,7 +938,7 @@ expand_teams_call (basic_block bb, gomp_teams *entry_stmt)
     num_teams = build_int_cst (unsigned_type_node, 0);
   else
     {
-      num_teams = OMP_CLAUSE_NUM_TEAMS_EXPR (num_teams);
+      num_teams = OMP_CLAUSE_NUM_TEAMS_UPPER_EXPR (num_teams);
       num_teams = fold_convert (unsigned_type_node, num_teams);
     }
   tree thread_limit = omp_find_clause (clauses, OMP_CLAUSE_THREAD_LIMIT);
@@ -9625,7 +9625,7 @@ get_target_arguments (gimple_stmt_iterator *gsi, gomp_target *tgt_stmt)
   tree clauses = gimple_omp_target_clauses (tgt_stmt);
   tree t, c = omp_find_clause (clauses, OMP_CLAUSE_NUM_TEAMS);
   if (c)
-    t = OMP_CLAUSE_NUM_TEAMS_EXPR (c);
+    t = OMP_CLAUSE_NUM_TEAMS_UPPER_EXPR (c);
   else
     t = integer_minus_one_node;
   push_target_argument_according_to_value (gsi, GOMP_TARGET_ARG_DEVICE_ALL,

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index f58a191e014cb751b4fc95129e47eb1b872792b3..d5841ea731320ad931dc6d10ab9a30f23df13013 100644 (file)
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -13906,7 +13906,7 @@ lower_omp_teams (gimple_stmt_iterator *gsi_p, omp_context *ctx)
     num_teams = build_int_cst (unsigned_type_node, 0);
   else
     {
-      num_teams = OMP_CLAUSE_NUM_TEAMS_EXPR (num_teams);
+      num_teams = OMP_CLAUSE_NUM_TEAMS_UPPER_EXPR (num_teams);
       num_teams = fold_convert (unsigned_type_node, num_teams);
       gimplify_expr (&num_teams, &bind_body, NULL, is_gimple_val, fb_rvalue);
     }

diff --git a/gcc/testsuite/c-c++-common/gomp/clauses-1.c b/gcc/testsuite/c-c++-common/gomp/clauses-1.c
index 742132f202eb873f962e0f605428dcec77dc0b8e..3ff49e0a29829ddab278804ccde7e3a972f8c637 100644 (file)
--- a/gcc/testsuite/c-c++-common/gomp/clauses-1.c
+++ b/gcc/testsuite/c-c++-common/gomp/clauses-1.c
@@ -164,7 +164,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
     ll++;
   #pragma omp target teams \
     device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \
-    shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) nowait depend(inout: dd[0]) \
+    shared(s) default(shared) reduction(+:r) num_teams(nte - 1:nte) thread_limit(tl) nowait depend(inout: dd[0]) \
     allocate (omp_default_mem_alloc:f) in_reduction(+:r2)
     ;
   #pragma omp target teams distribute \
@@ -175,7 +175,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
     ;
   #pragma omp target teams distribute parallel for \
     device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \
-    shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) \
+    shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl) \
     collapse(1) dist_schedule(static, 16) \
     if (parallel: i2) num_threads (nth) proc_bind(spread) \
     lastprivate (l) schedule(static, 4) nowait depend(inout: dd[0]) order(concurrent) \
@@ -194,7 +194,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
     ll++;
   #pragma omp target teams distribute simd \
     device(d) map (tofrom: m) if (i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \
-    shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) \
+    shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl) \
     collapse(1) dist_schedule(static, 16) order(concurrent) \
     safelen(8) simdlen(4) aligned(q: 32) nowait depend(inout: dd[0]) nontemporal(ntm) \
     allocate (omp_default_mem_alloc:f) in_reduction(+:r2)
@@ -236,7 +236,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
     ;
   #pragma omp target
   #pragma omp teams distribute parallel for \
-    private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) \
+    private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl) \
     collapse(1) dist_schedule(static, 16) \
     if (parallel: i2) num_threads (nth) proc_bind(spread) \
     lastprivate (l) schedule(static, 4) order(concurrent) allocate (omp_default_mem_alloc: f)
@@ -254,7 +254,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
     ll++;
   #pragma omp target
   #pragma omp teams distribute simd \
-    private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) \
+    private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl) \
     collapse(1) dist_schedule(static, 16) order(concurrent) \
     safelen(8) simdlen(4) aligned(q: 32) if(i3) nontemporal(ntm) \
     allocate (omp_default_mem_alloc: f)
@@ -268,7 +268,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
   for (int i = 0; i < 64; i++)
     ll++;
   #pragma omp teams distribute parallel for \
-    private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) \
+    private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl) \
     collapse(1) dist_schedule(static, 16) order(concurrent) \
     if (parallel: i2) num_threads (nth) proc_bind(spread) \
     lastprivate (l) schedule(static, 4) allocate (f)
@@ -284,7 +284,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
   for (int i = 0; i < 64; i++)
     ll++;
   #pragma omp teams distribute parallel for simd \
-    private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) \
+    private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl) \
     collapse(1) dist_schedule(static, 16) \
     if (parallel: i2) num_threads (nth) proc_bind(spread) \
     lastprivate (l) schedule(static, 4) order(concurrent) \
@@ -417,7 +417,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
   for (l = 0; l < 64; l++)
     ll++;
   #pragma omp teams loop \
-    private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) \
+    private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl) \
     collapse(1) lastprivate (l) bind(teams) allocate (f)
   for (l = 0; l < 64; ++l)
     ;
@@ -442,7 +442,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
     ;
   #pragma omp target teams loop \
     device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \
-    shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) nowait depend(inout: dd[0]) \
+    shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl) nowait depend(inout: dd[0]) \
     lastprivate (l) bind(teams) collapse(1) \
     allocate (omp_default_mem_alloc: f) in_reduction(+:r2)
   for (l = 0; l < 64; ++l)

diff --git a/gcc/testsuite/c-c++-common/gomp/num-teams-1.c b/gcc/testsuite/c-c++-common/gomp/num-teams-1.c
new file mode 100644 (file)
index 0000000..50cad85
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/num-teams-1.c
@@ -0,0 +1,48 @@
+int fn (int);
+
+void
+foo (void)
+{
+  #pragma omp teams num_teams (4 : 6)
+  ;
+  #pragma omp teams num_teams (7)
+  ;
+}
+
+void
+bar (void)
+{
+  #pragma omp target teams num_teams (5 : 19)
+  ;
+  #pragma omp target teams num_teams (21)
+  ;
+}
+
+void
+baz (void)
+{
+  #pragma omp teams num_teams (fn (1) : fn (2))
+  ;
+  #pragma omp teams num_teams (fn (3))
+  ;
+}
+
+void
+qux (void)
+{
+  #pragma omp target teams num_teams (fn (4) : fn (5))
+  ;
+  #pragma omp target teams num_teams (fn (6))
+  ;
+}
+
+void
+corge (void)
+{
+  #pragma omp target
+  #pragma omp teams num_teams (fn (7) : fn (8))
+  ;
+  #pragma omp target
+  #pragma omp teams num_teams (fn (9))
+  ;
+}

diff --git a/gcc/testsuite/c-c++-common/gomp/num-teams-2.c b/gcc/testsuite/c-c++-common/gomp/num-teams-2.c
new file mode 100644 (file)
index 0000000..242b994
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/num-teams-2.c
@@ -0,0 +1,27 @@
+int fn (int);
+
+void
+foo (int i)
+{
+  #pragma omp teams num_teams (6 : 4)          /* { dg-warning "'num_teams' lower bound '6' bigger than upper bound '4'" } */
+  ;
+  #pragma omp teams num_teams (-7)             /* { dg-warning "'num_teams' value must be positive" } */
+  ;
+  #pragma omp teams num_teams (i : -7)         /* { dg-warning "'num_teams' value must be positive" } */
+  ;
+  #pragma omp teams num_teams (-7 : 8)         /* { dg-warning "'num_teams' value must be positive" } */
+  ;
+}
+
+void
+bar (int i)
+{
+  #pragma omp target teams num_teams (6 : 4)   /* { dg-warning "'num_teams' lower bound '6' bigger than upper bound '4'" } */
+  ;
+  #pragma omp target teams num_teams (-7)      /* { dg-warning "'num_teams' value must be positive" } */
+  ;
+  #pragma omp target teams num_teams (i : -7)  /* { dg-warning "'num_teams' value must be positive" } */
+  ;
+  #pragma omp target teams num_teams (-7 : 8)  /* { dg-warning "'num_teams' value must be positive" } */
+  ;
+}

diff --git a/gcc/testsuite/g++.dg/gomp/attrs-1.C b/gcc/testsuite/g++.dg/gomp/attrs-1.C
index 2a5f2cf63234415b513c12d4839074311ba9db37..319ad3241de7dc99e8f4fd5fe76598ebd14889ed 100644 (file)
--- a/gcc/testsuite/g++.dg/gomp/attrs-1.C
+++ b/gcc/testsuite/g++.dg/gomp/attrs-1.C
@@ -211,7 +211,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
     ll++;
   [[omp::sequence (directive (target teams
     device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp)
-    shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) nowait depend(inout: dd[0])
+    shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl) nowait depend(inout: dd[0])
     allocate (omp_default_mem_alloc:f) in_reduction(+:r2)))]]
     ;
   [[omp::sequence (directive (target
@@ -226,7 +226,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
     ;
   [[omp::directive (target teams distribute parallel for
     device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp)
-    shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl)
+    shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl)
     collapse(1) dist_schedule(static, 16)
     if (parallel: i2) num_threads (nth) proc_bind(spread)
     lastprivate (l) schedule(static, 4) nowait depend(inout: dd[0]) order(concurrent)
@@ -245,7 +245,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
     ll++;
   [[omp::directive (target teams distribute simd
     device(d) map (tofrom: m) if (i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp)
-    shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl)
+    shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl)
     collapse(1) dist_schedule(static, 16) order(concurrent)
     safelen(8) simdlen(4) aligned(q: 32) nowait depend(inout: dd[0]) nontemporal(ntm)
     allocate (omp_default_mem_alloc:f) in_reduction(+:r2))]]
@@ -309,7 +309,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
   for (int i = 0; i < 64; i++)
     ;
   [[omp::directive (teams
-    private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl)
+    private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl)
     allocate (omp_default_mem_alloc: f))]]
     ;
   [[omp::sequence (omp::directive (target),
@@ -322,7 +322,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
     ll++;
   [[omp::sequence (directive (target),
     directive (teams distribute parallel for simd
-    private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl)
+    private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl)
     collapse(1) dist_schedule(static, 16)
     if (parallel: i2) num_threads (nth) proc_bind(spread)
     lastprivate (l) schedule(static, 4) order(concurrent)
@@ -339,7 +339,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
   for (int i = 0; i < 64; i++)
     ll++;
   [[omp::directive (teams distribute parallel for
-    private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl)
+    private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl)
     collapse(1) dist_schedule(static, 16)
     if (parallel: i2) num_threads (nth) proc_bind(spread)
     lastprivate (l) schedule(static, 4) copyin(t) allocate (f))]]
@@ -353,7 +353,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
   for (int i = 0; i < 64; i++)
     ll++;
   [[omp::directive (teams distribute parallel for simd
-    private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl)
+    private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl)
     collapse(1) dist_schedule(static, 16)
     if (parallel: i2) num_threads (nth) proc_bind(spread)
     lastprivate (l) schedule(static, 4)
@@ -371,7 +371,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
   for (int i = 0; i < 64; i++)
     ll++;
   [[omp::directive (teams distribute simd
-    private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl)
+    private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl)
     collapse(1) dist_schedule(static, 16) order(concurrent)
     safelen(8) simdlen(4) aligned(q: 32) if(i3) nontemporal(ntm) allocate(f))]]
   for (int i = 0; i < 64; i++)
@@ -507,7 +507,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
   for (l = 0; l < 64; ++l)
     ;
   [[omp::directive (teams loop
-    private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl)
+    private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte - 1:nte) thread_limit(tl)
     collapse(1) lastprivate (l) order(concurrent) allocate (f))]]
   for (l = 0; l < 64; ++l)
     ;
@@ -534,7 +534,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
     ;
   [[omp::directive (target teams loop
     device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp)
-    shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) nowait depend(inout: dd[0])
+    shared(s) default(shared) reduction(+:r) num_teams(nte - 1 : nte) thread_limit(tl) nowait depend(inout: dd[0])
     lastprivate (l) order(concurrent) collapse(1)
     allocate (omp_default_mem_alloc: f) in_reduction(+:r2))]]
   for (l = 0; l < 64; ++l)

diff --git a/gcc/testsuite/g++.dg/gomp/attrs-2.C b/gcc/testsuite/g++.dg/gomp/attrs-2.C
index c00be7f1db71e46bac71cfb2b642e9190dbdd3d8..955b2dd04c7b8805dcc4e694aef53a727cec97ef 100644 (file)
--- a/gcc/testsuite/g++.dg/gomp/attrs-2.C
+++ b/gcc/testsuite/g++.dg/gomp/attrs-2.C
@@ -220,7 +220,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
     ;
   [[omp::sequence (omp::directive (target teams distribute,
     device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp),
-    shared(s),default(shared),reduction(+:r),num_teams(nte),thread_limit(tl),order(concurrent),
+    shared(s),default(shared),reduction(+:r),num_teams(nte-1:nte),thread_limit(tl),order(concurrent),
     collapse(1),dist_schedule(static, 16),nowait depend(inout: dd[0]),allocate (omp_default_mem_alloc:f),in_reduction(+:r2)))]]
   for (int i = 0; i < 64; i++)
     ;
@@ -235,7 +235,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
     ll++;
   [[omp::directive (target teams distribute parallel for simd,
     device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp),
-    shared(s),default(shared),reduction(+:r),num_teams(nte),thread_limit(tl),
+    shared(s),default(shared),reduction(+:r),num_teams(2:nte),thread_limit(tl),
     collapse(1),dist_schedule(static, 16),
     if (parallel: i2),num_threads (nth),proc_bind(spread),
     lastprivate (l),schedule(static, 4),order(concurrent),
@@ -304,7 +304,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
   [[omp::directive (taskwait)]];
   [[omp::sequence (directive (target, nowait,depend(inout: dd[0]),in_reduction(+:r2)),
     directive (teams distribute,
-    private(p),firstprivate (f),shared(s),default(shared),reduction(+:r),num_teams(nte),thread_limit(tl),
+    private(p),firstprivate(f),shared(s),default(shared),reduction(+:r),num_teams(nte - 1 : nte),thread_limit(tl),
     collapse(1),dist_schedule(static, 16),allocate (omp_default_mem_alloc: f),order(concurrent)))]]
   for (int i = 0; i < 64; i++)
     ;
@@ -314,7 +314,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
     ;
   [[omp::sequence (omp::directive (target),
     omp::directive (teams distribute parallel for,
-    private(p),firstprivate (f),shared(s),default(shared),reduction(+:r),num_teams(nte),thread_limit(tl),
+    private(p),firstprivate(f),shared(s),default(shared),reduction(+:r),num_teams(16:nte),thread_limit(tl),
     collapse(1),dist_schedule(static, 16),
     if (parallel: i2),num_threads (nth),proc_bind(spread),
     lastprivate (l),schedule(static, 4),order(concurrent),allocate (omp_default_mem_alloc: f)))]]
@@ -332,7 +332,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
     ll++;
   [[omp::sequence (directive (target),
     directive (teams distribute simd,
-    private(p),firstprivate (f),shared(s),default(shared),reduction(+:r),num_teams(nte),thread_limit(tl),
+    private(p),firstprivate(f),shared(s),default(shared),reduction(+:r),num_teams(nte-1:nte),thread_limit(tl),
     collapse(1),dist_schedule(static, 16),order(concurrent),
     safelen(8),simdlen(4),aligned(q: 32),if(i3),nontemporal(ntm),
     allocate (omp_default_mem_alloc: f)))]]
@@ -346,7 +346,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
   for (int i = 0; i < 64; i++)
     ll++;
   [[omp::directive (teams distribute parallel for,
-    private(p),firstprivate (f),shared(s),default(shared),reduction(+:r),num_teams(nte),thread_limit(tl),
+    private(p),firstprivate (f),shared(s),default(shared),reduction(+:r),num_teams(nte-1:nte),thread_limit(tl),
     collapse(1),dist_schedule(static, 16),order(concurrent),
     if (parallel: i2),num_threads (nth),proc_bind(spread),
     lastprivate (l),schedule(static, 4),allocate (f))]]
@@ -362,7 +362,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
   for (int i = 0; i < 64; i++)
     ll++;
   [[omp::directive (teams distribute parallel for simd,
-    private(p),firstprivate (f),shared(s),default(shared),reduction(+:r),num_teams(nte),thread_limit(tl),
+    private(p),firstprivate (f),shared(s),default(shared),reduction(+:r),num_teams(nte-1:nte),thread_limit(tl),
     collapse(1),dist_schedule(static, 16),
     if (parallel: i2),num_threads (nth),proc_bind(spread),
     lastprivate (l),schedule(static, 4),order(concurrent),
@@ -502,7 +502,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
   for (l = 0; l < 64; l++)
     ll++;
   [[omp::directive (teams loop,
-    private(p),firstprivate (f),shared(s),default(shared),reduction(+:r),num_teams(nte),thread_limit(tl),
+    private(p),firstprivate (f),shared(s),default(shared),reduction(+:r),num_teams(nte-1:nte),thread_limit(tl),
     collapse(1),lastprivate (l),bind(teams),allocate (f))]]
   for (l = 0; l < 64; ++l)
     ;
@@ -527,7 +527,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
     ;
   [[omp::directive (target teams loop,
     device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp),
-    shared(s),default(shared),reduction(+:r),num_teams(nte),thread_limit(tl),nowait,depend(inout: dd[0]),
+    shared(s),default(shared),reduction(+:r),num_teams(nte-1:nte),thread_limit(tl),nowait,depend(inout: dd[0]),
     lastprivate (l),bind(teams),collapse(1),
     allocate (omp_default_mem_alloc: f),in_reduction(+:r2))]]
   for (l = 0; l < 64; ++l)

diff --git a/gcc/testsuite/g++.dg/gomp/num-teams-1.C b/gcc/testsuite/g++.dg/gomp/num-teams-1.C
new file mode 100644 (file)
index 0000000..5b36ffb
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/num-teams-1.C
@@ -0,0 +1,122 @@
+int fn1 (int);
+template <typename T>
+T fn2 (T);
+
+template <int N>
+void
+f1 ()
+{
+  #pragma omp teams num_teams (4 : 6)
+  ;
+  #pragma omp teams num_teams (7)
+  ;
+}
+
+template <int N>
+void
+f2 ()
+{
+  #pragma omp target teams num_teams (5 : 19)
+  ;
+  #pragma omp target teams num_teams (21)
+  ;
+}
+
+template <int N>
+void
+f3 ()
+{
+  #pragma omp teams num_teams (fn1 (1) : fn1 (2))
+  ;
+  #pragma omp teams num_teams (fn1 (3))
+  ;
+}
+
+template <int N>
+void
+f4 ()
+{
+  #pragma omp target teams num_teams (fn1 (4) : fn1 (5))
+  ;
+  #pragma omp target teams num_teams (fn1 (6))
+  ;
+}
+
+template <int N>
+void
+f5 ()
+{
+  #pragma omp target
+  #pragma omp teams num_teams (fn1 (7) : fn1 (8))
+  ;
+  #pragma omp target
+  #pragma omp teams num_teams (fn1 (9))
+  ;
+}
+
+template <typename T, T N4, T N6, T N7>
+void
+f1 ()
+{
+  #pragma omp teams num_teams (N4 : N6)
+  ;
+  #pragma omp teams num_teams (N7)
+  ;
+}
+
+template <typename T, T N5, T N19, T N21>
+void
+f2 ()
+{
+  #pragma omp target teams num_teams (N5 : N19)
+  ;
+  #pragma omp target teams num_teams (N21)
+  ;
+}
+
+template <typename T, T N1, T N2, T N3>
+void
+f3 ()
+{
+  #pragma omp teams num_teams (fn2 (N1) : fn2 (N2))
+  ;
+  #pragma omp teams num_teams (fn2 (N3))
+  ;
+}
+
+template <typename T, T N4, T N5, T N6>
+void
+f4 ()
+{
+  #pragma omp target teams num_teams (fn2 (N4) : fn2 (N5))
+  ;
+  #pragma omp target teams num_teams (fn2 (N6))
+  ;
+}
+
+template <typename T, T N7, T N8, T N9>
+void
+f5 ()
+{
+  #pragma omp target
+  #pragma omp teams num_teams (fn2 (N7) : fn2 (N8))
+  ;
+  #pragma omp target
+  #pragma omp teams num_teams (fn2 (N9))
+  ;
+}
+
+void
+test ()
+{
+  f1<0> ();
+  f2<0> ();
+  f3<0> ();
+  f4<0> ();
+  f5<0> ();
+  f1<int, 4, 6, 7> ();
+  f2<int, 5, 19, 21> ();
+  f3<int, 1, 2, 3> ();
+  f4<int, 4, 5, 6> ();
+  f5<int, 7, 8, 9> ();
+}

diff --git a/gcc/testsuite/g++.dg/gomp/num-teams-2.C b/gcc/testsuite/g++.dg/gomp/num-teams-2.C
new file mode 100644 (file)
index 0000000..8b8933c
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/num-teams-2.C
@@ -0,0 +1,64 @@
+template <int N>
+void
+foo (int i)
+{
+  #pragma omp teams num_teams (6 : 4)          // { dg-warning "'num_teams' lower bound '6' bigger than upper bound '4'" }
+  ;
+  #pragma omp teams num_teams (-7)             // { dg-warning "'num_teams' value must be positive" }
+  ;
+  #pragma omp teams num_teams (i : -7)         // { dg-warning "'num_teams' value must be positive" }
+  ;
+  #pragma omp teams num_teams (-7 : 8)         // { dg-warning "'num_teams' value must be positive" }
+  ;
+}
+
+template <int N>
+void
+bar (int i)
+{
+  #pragma omp target teams num_teams (6 : 4)   // { dg-warning "'num_teams' lower bound '6' bigger than upper bound '4'" }
+  ;
+  #pragma omp target teams num_teams (-7)      // { dg-warning "'num_teams' value must be positive" }
+  ;
+  #pragma omp target teams num_teams (i : -7)  // { dg-warning "'num_teams' value must be positive" }
+  ;
+  #pragma omp target teams num_teams (-7 : 8)  // { dg-warning "'num_teams' value must be positive" }
+  ;
+}
+
+template <typename T, T NM7, T N4, T N6, T N8>
+void
+baz (T i)
+{
+  #pragma omp teams num_teams (N6 : N4)                // { dg-warning "'num_teams' lower bound '6' bigger than upper bound '4'" }
+  ;
+  #pragma omp teams num_teams (NM7)            // { dg-warning "'num_teams' value must be positive" }
+  ;
+  #pragma omp teams num_teams (i : NM7)                // { dg-warning "'num_teams' value must be positive" }
+  ;
+  #pragma omp teams num_teams (NM7 : N8)       // { dg-warning "'num_teams' value must be positive" }
+  ;
+}
+
+template <typename T, T NM7, T N4, T N6, T N8>
+void
+qux (T i)
+{
+  #pragma omp target teams num_teams (N6 : N4) // { dg-warning "'num_teams' lower bound '6' bigger than upper bound '4'" }
+  ;
+  #pragma omp target teams num_teams (NM7)     // { dg-warning "'num_teams' value must be positive" }
+  ;
+  #pragma omp target teams num_teams (i : NM7) // { dg-warning "'num_teams' value must be positive" }
+  ;
+  #pragma omp target teams num_teams (NM7 : N8)        // { dg-warning "'num_teams' value must be positive" }
+  ;
+}
+
+void
+test ()
+{
+  foo<0> (5);
+  bar<0> (5);
+  baz<int, -7, 4, 6, 8> (5);
+  qux<int, -7, 4, 6, 8> (5);
+}

diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 275dc7d8af737ab620dd576c7c77501c26ca8018..f6383b91efae8626a402e0e048f54d3d39835c19 100644 (file)
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -994,7 +994,13 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
 
     case OMP_CLAUSE_NUM_TEAMS:
       pp_string (pp, "num_teams(");
-      dump_generic_node (pp, OMP_CLAUSE_NUM_TEAMS_EXPR (clause),
+      if (OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR (clause))
+       {
+         dump_generic_node (pp, OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR (clause),
+                            spc, flags, false);
+         pp_colon (pp);
+       }
+      dump_generic_node (pp, OMP_CLAUSE_NUM_TEAMS_UPPER_EXPR (clause),
                         spc, flags, false);
       pp_right_paren (pp);
       break;

diff --git a/gcc/tree.c b/gcc/tree.c
index 845228a055b2cfac0c9ca8c0cda1b9df4b0095c6..f2c829fa4c61cf8f2d89da7d3a1bce55f780e423 100644 (file)
--- a/gcc/tree.c
+++ b/gcc/tree.c
@@ -330,7 +330,7 @@ unsigned const char omp_clause_num_ops[] =
   1, /* OMP_CLAUSE_DIST_SCHEDULE  */
   0, /* OMP_CLAUSE_INBRANCH  */
   0, /* OMP_CLAUSE_NOTINBRANCH  */
-  1, /* OMP_CLAUSE_NUM_TEAMS  */
+  2, /* OMP_CLAUSE_NUM_TEAMS  */
   1, /* OMP_CLAUSE_THREAD_LIMIT  */
   0, /* OMP_CLAUSE_PROC_BIND  */
   1, /* OMP_CLAUSE_SAFELEN  */

diff --git a/gcc/tree.h b/gcc/tree.h
index f62c00bc8707029db52e2f3fe529948755235d3d..92c3d77f09f0dd0522c20233408c3f20d0d276d7 100644 (file)
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1803,9 +1803,12 @@ class auto_suppress_location_wrappers
 #define OMP_CLAUSE_ALLOCATE_COMBINED(NODE) \
   (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_ALLOCATE)->base.public_flag)
 
-#define OMP_CLAUSE_NUM_TEAMS_EXPR(NODE) \
+#define OMP_CLAUSE_NUM_TEAMS_UPPER_EXPR(NODE) \
   OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_NUM_TEAMS), 0)
 
+#define OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR(NODE) \
+  OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_NUM_TEAMS), 1)
+
 #define OMP_CLAUSE_THREAD_LIMIT_EXPR(NODE) \
   OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, \
                                                OMP_CLAUSE_THREAD_LIMIT), 0)

diff --git a/libgomp/testsuite/libgomp.c-c++-common/teams-1.c b/libgomp/testsuite/libgomp.c-c++-common/teams-1.c
new file mode 100644 (file)
index 0000000..76189ef
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/teams-1.c
@@ -0,0 +1,26 @@
+#include <omp.h>
+#include <stdlib.h>
+
+int
+main ()
+{
+  #pragma omp teams num_teams (5)
+  {
+    if (omp_get_num_teams () != 5)
+      abort ();
+    #pragma omp distribute dist_schedule(static,1)
+    for (int i = 0; i < 5; ++i)
+      if (omp_get_team_num () != i)
+       abort ();
+  }
+  #pragma omp teams num_teams (7 : 9)
+  {
+    if (omp_get_num_teams () < 7 || omp_get_num_teams () > 9)
+      abort ();
+    #pragma omp distribute dist_schedule(static,1)
+    for (int i = 0; i < omp_get_num_teams (); ++i)
+      if (omp_get_team_num () != i)
+       abort ();
+  }
+  return 0;
+}