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.