Re: [committed] openmp: Add support for 2 argument num_teams clause

["H.J. Lu" <hjl.tools@gmail.com>]

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  <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.
>

This caused:

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=103224

May need bootstrap to reproduce it.

-- 
H.J.