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.
(cherry picked from commit
48d7327f2aaf65e224f5f0793a65b950297f6c7f)
+2022-02-27 Tobias Burnus <tobias@codesourcery.com>
+
+ Backported from master:
+ 2021-11-11 Jakub Jelinek <jakub@redhat.com>
+
+ * 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.
+
2022-02-27 Tobias Burnus <tobias@codesourcery.com>
Backported from master:
+2022-02-27 Tobias Burnus <tobias@codesourcery.com>
+
+ Backported from master:
+ 2021-11-11 Jakub Jelinek <jakub@redhat.com>
+
+ * 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.
+
2022-02-18 Kwok Cheung Yeung <kcy@codesourcery.com>
* c-parser.c (c_parser_omp_construct): Move handling of
}
/* 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)
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;
}
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];
+2022-02-27 Tobias Burnus <tobias@codesourcery.com>
+
+ Backported from master:
+ 2021-11-11 Jakub Jelinek <jakub@redhat.com>
+
+ * 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.
+
2022-02-18 Kwok Cheung Yeung <kcy@codesourcery.com>
* parser.c (cp_parser_omp_construct): Move handling of
}
/* 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,
"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;
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;
+ }
c_omp_adjust_map_clauses (cclauses[C_OMP_CLAUSE_SPLIT_TARGET], true);
finish_omp_target (pragma_tok->location,
cclauses[C_OMP_CLAUSE_SPLIT_TARGET], body, true);
= 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:
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:
{
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;
}
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:
+2022-02-27 Tobias Burnus <tobias@codesourcery.com>
+
+ Backported from master:
+ 2021-11-11 Jakub Jelinek <jakub@redhat.com>
+
+ * trans-openmp.c (gfc_trans_omp_clauses): Use
+ OMP_CLAUSE_NUM_TEAMS_UPPER_EXPR instead of OMP_CLAUSE_NUM_TEAMS_EXPR.
+
2022-02-11 Kwok Cheung Yeung <kcy@codesourcery.com>
* gfortran.h (is_omp_declarative_stmt): New.
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);
}
= 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:
{
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);
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)
{
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;
}
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);
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,
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);
}
+2022-02-27 Tobias Burnus <tobias@codesourcery.com>
+
+ Backported from master:
+ 2021-11-11 Jakub Jelinek <jakub@redhat.com>
+
+ * 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.
+
2022-02-27 Tobias Burnus <tobias@codesourcery.com>
Backported from master:
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 \
;
#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) \
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)
;
#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)
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)
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)
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) \
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)
;
;
#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)
--- /dev/null
+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))
+ ;
+}
--- /dev/null
+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" } */
+ ;
+}
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
;
[[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)
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))]]
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),
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)
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))]]
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)
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++)
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)
;
;
[[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)
;
[[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++)
;
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),
[[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++)
;
;
[[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)))]]
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)))]]
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))]]
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),
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)
;
;
[[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)
--- /dev/null
+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> ();
+}
--- /dev/null
+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);
+}
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;
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 */
#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)
+2022-02-27 Tobias Burnus <tobias@codesourcery.com>
+
+ Backported from master:
+ 2021-11-11 Jakub Jelinek <jakub@redhat.com>
+
+ * testsuite/libgomp.c-c++-common/teams-1.c: New test.
+
2022-02-27 Tobias Burnus <tobias@codesourcery.com>
Backported from master:
--- /dev/null
+#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;
+}
projects / thirdparty / gcc.git / commitdiff
