aboutsummaryrefslogtreecommitdiff
path: root/gcc
diff options
context:
space:
mode:
Diffstat (limited to 'gcc')
-rw-r--r--gcc/c/c-parser.c111
-rw-r--r--gcc/cp/parser.c78
-rw-r--r--gcc/cp/pt.c59
-rw-r--r--gcc/cp/semantics.c47
-rw-r--r--gcc/fortran/trans-openmp.c2
-rw-r--r--gcc/gimplify.c55
-rw-r--r--gcc/omp-expand.c4
-rw-r--r--gcc/omp-low.c2
-rw-r--r--gcc/testsuite/c-c++-common/gomp/clauses-1.c18
-rw-r--r--gcc/testsuite/c-c++-common/gomp/num-teams-1.c48
-rw-r--r--gcc/testsuite/c-c++-common/gomp/num-teams-2.c27
-rw-r--r--gcc/testsuite/g++.dg/gomp/attrs-1.C20
-rw-r--r--gcc/testsuite/g++.dg/gomp/attrs-2.C18
-rw-r--r--gcc/testsuite/g++.dg/gomp/num-teams-1.C122
-rw-r--r--gcc/testsuite/g++.dg/gomp/num-teams-2.C64
-rw-r--r--gcc/tree-pretty-print.c8
-rw-r--r--gcc/tree.c2
-rw-r--r--gcc/tree.h5
18 files changed, 556 insertions, 134 deletions
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 80dd61d..3c9f587 100644
--- 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 32de97b..4c1fed0 100644
--- 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 b2916f8..82bf7dc 100644
--- 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 2443d03..60e0982 100644
--- 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 e81c558..22d6662 100644
--- 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 c2ab96e..e5877c83 100644
--- 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 70957a6..1a5426e 100644
--- 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 f58a191..d5841ea 100644
--- 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 742132f..3ff49e0 100644
--- 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
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
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 2a5f2cf..319ad32 100644
--- 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 c00be7f..955b2dd 100644
--- 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
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
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 275dc7d..f6383b9 100644
--- 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 845228a..f2c829f 100644
--- 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 f62c00b..92c3d77 100644
--- 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)
generated by cgit v1.1 at 2025-05-28 20:01:10 +0000