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

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.

	Jakub

Message ID	20211111091130.GN2710@tucnak
State	New
Headers	show Return-Path: <gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org> DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org F03C7385843E Date: Thu, 11 Nov 2021 10:11:30 +0100 To: gcc-patches@gcc.gnu.org Subject: [committed] openmp: Add support for 2 argument num_teams clause Message-ID: <20211111091130.GN2710@tucnak> MIME-Version: 1.0 Content-Type: text/plain; charset=us-ascii Content-Disposition: inline Precedence: list From: Jakub Jelinek via Gcc-patches <gcc-patches@gcc.gnu.org> Reply-To: Jakub Jelinek <jakub@redhat.com> Cc: Tobias Burnus <tobias@codesourcery.com> Errors-To: gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org Sender: "Gcc-patches" <gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org>
Series	[committed] openmp: Add support for 2 argument num_teams clause \| expand [committed] openmp: Add support for 2 argument num_teams clause

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

Commit Message

Comments

Patch