From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1534) id 689783858435; Sun, 27 Feb 2022 21:34:59 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 689783858435 Content-Type: text/plain; charset="us-ascii" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit From: Tobias Burnus To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/omp/gcc-11] openmp: Add support for 2 argument num_teams clause X-Act-Checkin: gcc X-Git-Author: Jakub Jelinek X-Git-Refname: refs/heads/devel/omp/gcc-11 X-Git-Oldrev: b8cd48d78904d51b1a1b7594f6c9307843fc0d73 X-Git-Newrev: a2662a871fe17bceb839d191fd8840dc2aba7e22 Message-Id: <20220227213459.689783858435@sourceware.org> Date: Sun, 27 Feb 2022 21:34:59 +0000 (GMT) X-BeenThere: gcc-cvs@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-cvs mailing list List-Unsubscribe: , List-Archive: List-Help: List-Subscribe: , X-List-Received-Date: Sun, 27 Feb 2022 21:34:59 -0000 https://gcc.gnu.org/g:a2662a871fe17bceb839d191fd8840dc2aba7e22 commit a2662a871fe17bceb839d191fd8840dc2aba7e22 Author: Jakub Jelinek Date: Thu Nov 11 09:42:47 2021 +0100 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 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) Diff: --- gcc/ChangeLog.omp | 20 ++++ gcc/c/ChangeLog.omp | 12 +++ gcc/c/c-parser.c | 111 ++++++++++++++------- gcc/cp/ChangeLog.omp | 17 ++++ gcc/cp/parser.c | 78 +++++++++------ gcc/cp/pt.c | 59 ++++++----- gcc/cp/semantics.c | 47 +++++++++ gcc/fortran/ChangeLog.omp | 8 ++ gcc/fortran/trans-openmp.c | 2 +- gcc/gimplify.c | 55 +++++++++- gcc/omp-expand.c | 4 +- gcc/omp-low.c | 2 +- gcc/testsuite/ChangeLog.omp | 15 +++ gcc/testsuite/c-c++-common/gomp/clauses-1.c | 18 ++-- gcc/testsuite/c-c++-common/gomp/num-teams-1.c | 48 +++++++++ gcc/testsuite/c-c++-common/gomp/num-teams-2.c | 27 +++++ gcc/testsuite/g++.dg/gomp/attrs-1.C | 20 ++-- gcc/testsuite/g++.dg/gomp/attrs-2.C | 18 ++-- gcc/testsuite/g++.dg/gomp/num-teams-1.C | 122 +++++++++++++++++++++++ gcc/testsuite/g++.dg/gomp/num-teams-2.C | 64 ++++++++++++ gcc/tree-pretty-print.c | 8 +- gcc/tree.c | 2 +- gcc/tree.h | 5 +- libgomp/ChangeLog.omp | 7 ++ libgomp/testsuite/libgomp.c-c++-common/teams-1.c | 26 +++++ 25 files changed, 661 insertions(+), 134 deletions(-) diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index 6ae003c7600..33c521b3876 100644 --- a/gcc/ChangeLog.omp +++ b/gcc/ChangeLog.omp @@ -1,3 +1,23 @@ +2022-02-27 Tobias Burnus + + Backported from master: + 2021-11-11 Jakub Jelinek + + * 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 Backported from master: diff --git a/gcc/c/ChangeLog.omp b/gcc/c/ChangeLog.omp index c1f9e860d16..919a4733299 100644 --- a/gcc/c/ChangeLog.omp +++ b/gcc/c/ChangeLog.omp @@ -1,3 +1,15 @@ +2022-02-27 Tobias Burnus + + Backported from master: + 2021-11-11 Jakub Jelinek + + * 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 * c-parser.c (c_parser_omp_construct): Move handling of diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index bb463c6d8af..466435a62b2 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -15211,7 +15211,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) @@ -15220,34 +15223,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, "% value must be positive"); - t = integer_one_node; + warning_at (upper_loc, 0, "% 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, "% 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, "% 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; } @@ -21073,31 +21110,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/ChangeLog.omp b/gcc/cp/ChangeLog.omp index 03b72006567..0b041e834f7 100644 --- a/gcc/cp/ChangeLog.omp +++ b/gcc/cp/ChangeLog.omp @@ -1,3 +1,20 @@ +2022-02-27 Tobias Burnus + + Backported from master: + 2021-11-11 Jakub Jelinek + + * 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 * parser.c (cp_parser_omp_construct): Move handling of diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index c910aafd508..368bcb39f23 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -38037,21 +38037,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, @@ -38061,7 +38075,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; @@ -43878,32 +43893,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; + } 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); diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c index d33cd3d9473..6e40364a4e5 100644 --- a/gcc/cp/pt.c +++ b/gcc/cp/pt.c @@ -17460,6 +17460,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: @@ -17468,7 +17475,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: @@ -18969,31 +18975,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 87f996fec1f..a066873b7f0 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -7285,6 +7285,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, + "% 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/ChangeLog.omp b/gcc/fortran/ChangeLog.omp index e76884dbc0c..5ea95855449 100644 --- a/gcc/fortran/ChangeLog.omp +++ b/gcc/fortran/ChangeLog.omp @@ -1,3 +1,11 @@ +2022-02-27 Tobias Burnus + + Backported from master: + 2021-11-11 Jakub Jelinek + + * 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 * gfortran.h (is_omp_declarative_stmt): New. diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index 3bb4feb8fa5..4dd53f0bb12 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -3958,7 +3958,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 b0226e98d02..2af3c10e76b 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -10844,9 +10844,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: @@ -14358,7 +14373,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); @@ -14366,14 +14382,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) { @@ -14411,7 +14455,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 2f9f4cd1d48..5da8f2034c1 100644 --- a/gcc/omp-expand.c +++ b/gcc/omp-expand.c @@ -941,7 +941,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); @@ -9706,7 +9706,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 8f748a442c8..5a211e6555f 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -14628,7 +14628,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/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp index 02ab20e8faa..6fb3ebea05a 100644 --- a/gcc/testsuite/ChangeLog.omp +++ b/gcc/testsuite/ChangeLog.omp @@ -1,3 +1,18 @@ +2022-02-27 Tobias Burnus + + Backported from master: + 2021-11-11 Jakub Jelinek + + * 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 Backported from master: diff --git a/gcc/testsuite/c-c++-common/gomp/clauses-1.c b/gcc/testsuite/c-c++-common/gomp/clauses-1.c index 742132f202e..3ff49e0a298 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 00000000000..50cad856cb3 --- /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 00000000000..242b994da22 --- /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 2a5f2cf6323..319ad3241de 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 c00be7f1db7..955b2dd04c7 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 00000000000..5b36ffb91ed --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/num-teams-1.C @@ -0,0 +1,122 @@ +int fn1 (int); +template +T fn2 (T); + +template +void +f1 () +{ + #pragma omp teams num_teams (4 : 6) + ; + #pragma omp teams num_teams (7) + ; +} + +template +void +f2 () +{ + #pragma omp target teams num_teams (5 : 19) + ; + #pragma omp target teams num_teams (21) + ; +} + +template +void +f3 () +{ + #pragma omp teams num_teams (fn1 (1) : fn1 (2)) + ; + #pragma omp teams num_teams (fn1 (3)) + ; +} + +template +void +f4 () +{ + #pragma omp target teams num_teams (fn1 (4) : fn1 (5)) + ; + #pragma omp target teams num_teams (fn1 (6)) + ; +} + +template +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 +void +f1 () +{ + #pragma omp teams num_teams (N4 : N6) + ; + #pragma omp teams num_teams (N7) + ; +} + +template +void +f2 () +{ + #pragma omp target teams num_teams (N5 : N19) + ; + #pragma omp target teams num_teams (N21) + ; +} + +template +void +f3 () +{ + #pragma omp teams num_teams (fn2 (N1) : fn2 (N2)) + ; + #pragma omp teams num_teams (fn2 (N3)) + ; +} + +template +void +f4 () +{ + #pragma omp target teams num_teams (fn2 (N4) : fn2 (N5)) + ; + #pragma omp target teams num_teams (fn2 (N6)) + ; +} + +template +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 (); + f2 (); + f3 (); + f4 (); + f5 (); +} 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 00000000000..8b8933cd0d3 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/num-teams-2.C @@ -0,0 +1,64 @@ +template +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 +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 +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 +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 (5); + qux (5); +} diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c index 6eedca4f91f..f41587fa21d 100644 --- a/gcc/tree-pretty-print.c +++ b/gcc/tree-pretty-print.c @@ -1036,7 +1036,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 7320247e05f..9146d3e06de 100644 --- a/gcc/tree.c +++ b/gcc/tree.c @@ -331,7 +331,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 c93d1b34f9b..6b63e2e0f52 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1818,9 +1818,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/ChangeLog.omp b/libgomp/ChangeLog.omp index 0546effcbcd..e8304a09a69 100644 --- a/libgomp/ChangeLog.omp +++ b/libgomp/ChangeLog.omp @@ -1,3 +1,10 @@ +2022-02-27 Tobias Burnus + + Backported from master: + 2021-11-11 Jakub Jelinek + + * testsuite/libgomp.c-c++-common/teams-1.c: New test. + 2022-02-27 Tobias Burnus Backported from master: 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 index 00000000000..76189ef0330 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/teams-1.c @@ -0,0 +1,26 @@ +#include +#include + +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; +}