From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from mail-pl1-x633.google.com (mail-pl1-x633.google.com [IPv6:2607:f8b0:4864:20::633]) by sourceware.org (Postfix) with ESMTPS id 0D3933858413 for ; Sat, 13 Nov 2021 16:51:00 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 0D3933858413 Received: by mail-pl1-x633.google.com with SMTP id y8so5484783plg.1 for ; Sat, 13 Nov 2021 08:51:00 -0800 (PST) X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20210112; h=x-gm-message-state:mime-version:references:in-reply-to:from:date :message-id:subject:to:cc; bh=v6g93LYx0frscLmv7dPVG+UvPwF5tNug+4pXVu2tKro=; b=0m7wAFl5LY4G4WFCX+7bkHyOePNzn6ssIAnBIMNK4lr5pheKtn8qfbmhO+UH1v5GtY hUf0r3RfvQr0q/00o8QvVNICf1Y+/SgGo9ojug2P/7xtWXs5gh8URpReN21JfzDAgPds LD1ldMfHx4vyPfI/u+6ZmVEFj7mMizwVFfS0N1bP0fcLKwAb9HRKUoQNqVZIY32aapWq xBd8MdMWOsVyonJ76c58em+dDq3ZGsd6o86f6+41a1Lhr/AeLrLkM0qKA2XpTLAEEwEU 1MRIO1CKcdExY9fPh4JjDFGf+MKXmBENo4rUUPDaE6vJcKjnFU7eC1msaGeuEfhxFrKN CPuw== X-Gm-Message-State: AOAM530zZ/yVhlfta3oxUuSsczcy/8EE+Nw5o7pLos+7Af/IjzCQKDes 3lYNyBiToRYGgcL0c1uWLkddRGPeUJebOnEzbQo= X-Google-Smtp-Source: ABdhPJy2h9SkdJRPr0qo25O0TMPiHjsGeFvmY0UJoiL4d/hbrWw4lctmrlcEKerIMeDpMknBcoGJ4/5/PqOoODJGCEU= X-Received: by 2002:a17:90b:1e0e:: with SMTP id pg14mr46471006pjb.143.1636822258854; Sat, 13 Nov 2021 08:50:58 -0800 (PST) MIME-Version: 1.0 References: <20211111091130.GN2710@tucnak> In-Reply-To: <20211111091130.GN2710@tucnak> From: "H.J. Lu" Date: Sat, 13 Nov 2021 08:50:22 -0800 Message-ID: Subject: Re: [committed] openmp: Add support for 2 argument num_teams clause To: Jakub Jelinek Cc: GCC Patches , Tobias Burnus Content-Type: text/plain; charset="UTF-8" X-Spam-Status: No, score=-3023.1 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, DKIM_VALID_AU, DKIM_VALID_EF, FREEMAIL_FROM, KAM_SHORT, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.4 X-Spam-Checker-Version: SpamAssassin 3.4.4 (2020-01-24) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , X-List-Received-Date: Sat, 13 Nov 2021 16:51:02 -0000 On Thu, Nov 11, 2021 at 1:12 AM Jakub Jelinek via Gcc-patches 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 > > 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.