On Thu, Nov 11, 2021 at 1:12 AM Jakub Jelinek via Gcc-patches
<gcc-patches@gcc.gnu.org> wrote:
>
> Hi!
>
> 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.
>
> Bootstrapped/regtested on x86_64-linux and i686-linux, committed to trunk.
>
> 2021-11-11 Jakub Jelinek <ja...@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.
>
This caused:
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=103224
May need bootstrap to reproduce it.
--
H.J.
