From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id 2581C3A600F7 for ; Fri, 7 May 2021 12:36:13 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 2581C3A600F7 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Thomas_Schwinge@mentor.com IronPort-SDR: 4cdhOYR+3vOFqPhSHeOiXOaAdAK7bJuo+whlR5ByeKECzuHbSduhj883yakV2Ic/FKA8ldceIa 0/p5aiBWOtJ1tS5FcOAFL7hL7B2+vx4GPoFnLtBtN/bpMi7chjuEus/ClcL5+0aM4EgT0bRjMX HXtyf76ladzuO30fs7lqa19JyBCmPgFvxAMHPz+fBsZWK9pKmYyGjMX0BPfKpcfghBkiRlvDby rTYAAeMl9IPRdoYGhuFGOL20t+u703lOv8Vnlv4YuYRPO39hsQkvwZXJwctnE56IELhoAPOdZl raU= X-IronPort-AV: E=Sophos;i="5.82,280,1613462400"; d="scan'208,223";a="60957181" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 07 May 2021 04:36:11 -0800 IronPort-SDR: gDOCE88TsLk7en3iIMivxMwrtKj/64yH89dYb8ETVqPq7To36ZptWlmbBdMR5cuuHzzlwCdPkz N1i5UEWQcy5n6q9bhNXa1CSUBkjO25PsMkNMa/GyJg4yr5fzuUBzFuFm6MJVyp7L36ELlWTrgf /dt4C08zyOV+ZNwSm8zv0rqZ8NY5hsQ4d6nYrzeKas4kZdke0ixzT2UnGzpir2N6SwcARM/MmX mK8ny4mgEGM44viqIWnJQkfbZJ6mJRH1LFx4yc2fS99xL+Pp6FirKPAD63U4vZkcfi4GgURA9p be8= From: Thomas Schwinge To: Chung-Lin Tang CC: , Catherine Moore , Tobias Burnus Subject: Re: [PATCH, OG10, OpenMP 5.0, committed] Implement relaxation of implicit map vs. existing device mappings In-Reply-To: <09ae7f35-41fd-c83e-2054-18ad4433d5e3@gmail.com> References: <09ae7f35-41fd-c83e-2054-18ad4433d5e3@gmail.com> User-Agent: Notmuch/0.29.3+94~g74c3f1b (https://notmuchmail.org) Emacs/27.1 (x86_64-pc-linux-gnu) Date: Fri, 7 May 2021 14:35:59 +0200 Message-ID: <87tuneu3f4.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 Content-Type: multipart/mixed; boundary="=-=-=" X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-09.mgc.mentorg.com (139.181.222.9) To svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) X-Spam-Status: No, score=-12.5 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) 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: Fri, 07 May 2021 12:36:19 -0000 --=-=-= Content-Type: text/plain; charset="utf-8" Content-Transfer-Encoding: quoted-printable Hi Chung-Lin! On 2021-05-05T23:17:25+0800, Chung-Lin Tang via Gcc-patches wrote: > This patch implements relaxing the requirements when a map with the impli= cit attribute encounters > an overlapping existing map. [...] Oh, oh, these data mapping interfaces/semantics ares getting more and more "convoluted"... %-\ (Not your fault, of course.) Haven't looked in too much detail in the patch/implementation (I'm not very well-versend in the exact OpenMP semantics anyway), but I suppose we should do similar things for OpenACC, too. I think we even currently do have a gimplification-level "hack" to replicate data clauses' array bounds for implicit data clauses on compute constructs, if the default "complete" mapping is going to clash with a "limited" mapping that's specified in an outer OpenACC 'data' directive. (That, of course, doesn't work for the general case of non-lexical scoping, or dynamic OpenACC 'enter data', etc., I suppose) I suppose your method could easily replace and improve that; we shall look into that later. That said, in your patch, is this current implementation (explicitly) meant or not meant to be active for OpenACC, too, or just OpenMP (I couldn't quickly tell), and/or is it (implicitly?) a no-op for OpenACC? > As the OpenMP 5.0 spec describes on page 320, lines 18-27 (and 5.1 spec, > page 352, lines 13-22): > > "If a single contiguous part of the original storage of a list item with = an implicit data-mapping > attribute has corresponding storage in the device data environment prio= r to a task encountering the > construct that is associated with the map clause, only that part of the= original storage will have > corresponding storage in the device data environment as a result of the= map clause." > > Also tracked in the OpenMP spec context as issue #1463: > https://github.com/OpenMP/spec/issues/1463 > > The implementation inside the compiler is to of course, tag the implicitl= y created maps with some > indication of "implicit". I've done this with a OMP_CLAUSE_MAP_IMPLICIT_P= macro, using > 'base.deprecated_flag' underneath. > > There is an encoding of this as GOMP_MAP_IMPLICIT =3D=3D GOMP_MAP_FLAG_SP= ECIAL_3|GOMP_MAP_FLAG_SPECIAL_4 > in include/gomp-constants.h for the runtime, but I've intentionally avoid= ed exploding the entire > gimplify/omp-low with a new set of GOMP_MAP_IMPLICIT_TO/FROM/etc. symbols= , instead adding in the new > flag bits only at the final runtime call generation during omp-lowering. > > The rest is libgomp mapping taking care of the implicit case: allowing ma= p success if an existing > map is a proper subset of the new map, if the new map is implicit. Straig= htforward enough I think. Seems so -- based on my very quick look. ;-) > There are also some additions to print the implicit attribute during tree= pretty-printing, for that > reason some scan tests were updated. ACK, thanks. > Also, another adjustment in this patch is how implicitly created clauses = are added to the current > clause list in gimplify_adjust_omp_clauses(). Instead of simply appending= the new clauses to the end, > this patch adds them at the position "after initial non-map clauses, but = right before any existing > map clauses". Probably you haven't been testing such a configuration; I've just pushed "Fix up 'c-c++-common/goacc/firstprivate-mappings-1.c' for C, non-LP64" to devel/omp/gcc-10 branch in commit c51cc3b96f0b562deaffcfbcc51043aed216801a, see attached. > The reason for this is: when combined with other map clauses, for example= : > > #pragma omp target map(rec.ptr[:N]) > for (int i =3D 0; i < N; i++) > rec.ptr[i] +=3D 1; > > There will be an implicit map created for map(rec), because of the access= inside the target region. > The expectation is that 'rec' is implicitly mapped, and then the pointed = array-section part by 'rec.ptr' > will be mapped, and then attachment to the 'rec.ptr' field of the mapped = 'rec' (in that order). > > If the implicit 'map(rec)' is appended to the end, instead of placed befo= re other maps, the attachment > operation will not find anything to attach to, and the entire region will= fail. But that doesn't (negatively) affect user-visible semantics (OpenMP, and also OpenACC, if applicable), in that more/bigger objects then get mapped than were before? (I suppose not?) Please make sure to put any rationale (like you've posted above) into source code comments ('gcc/gimplify.c:gimplify_adjust_omp_clauses', I suppose), or even GCC internals (?) manual, if applicable. Gr=C3=BC=C3=9Fe Thomas > Note: this touches a bit on another issue which I will be sending a patch= for later: > per the discussion on omp-lang, an array section list item should *not* b= e mapping its base-pointer > (although an attachment attempt should exist), while in current GCC behav= ior, for struct member pointers > like 'rec.ptr' above, we do map it (which should be deemed incorrect). > > This means that as of right now, this modification of map order doesn't r= eally exhibit the above mentioned > behavior yet. I have included it as part of this patch because the "[impl= icit]" tree printing requires > modifying many gimple scan tests already, so including the test modificat= ions together seems more > manageable patch-wise. > > Tested with no regressions, and pushed to devel/omp/gcc-10. Will be submi= tting a mainline trunk version later. > > Chung-Lin > > 2021-05-05 Chung-Lin Tang > > include/ChangeLog: > > * gomp-constants.h (GOMP_MAP_IMPLICIT): New special map kind bits v= alue. > (GOMP_MAP_FLAG_SPECIAL_BITS): Define helper mask for whole set of > special map kind bits. > (GOMP_MAP_NONCONTIG_ARRAY_P): Adjust test for non-contiguous array = map > kind bits to be more specific. > (GOMP_MAP_IMPLICIT_P): New predicate macro for implicit map kinds. > > gcc/ChangeLog: > > * tree.h (OMP_CLAUSE_MAP_IMPLICIT_P): New access macro for 'implici= t' > bit, using 'base.deprecated_flag' field of tree_node. > * tree-pretty-print.c (dump_omp_clause): Add support for printing > implicit attribute in tree dumping. > * gimplify.c (gimplify_adjust_omp_clauses_1): > Set OMP_CLAUSE_MAP_IMPLICIT_P to 1 if map clause is implicitly crea= ted. > (gimplify_adjust_omp_clauses): Adjust place of adding implicitly cr= eated > clauses, from simple append, to starting of list, after non-map cla= uses. > * omp-low.c (lower_omp_target): Add GOMP_MAP_IMPLICIT bits into kin= d > values passed to libgomp for implicit maps. > > gcc/testsuite/ChangeLog: > > * c-c++-common/gomp/target-implicit-map-1.c: New test. > * c-c++-common/goacc/combined-reduction.c: Adjust scan test pattern= . > * c-c++-common/goacc/firstprivate-mappings-1.c: Likewise. > * c-c++-common/goacc/mdc-1.c: Likewise. > * c-c++-common/goacc/reduction-1.c: Likewise. > * c-c++-common/goacc/reduction-2.c: Likewise. > * c-c++-common/goacc/reduction-3.c: Likewise. > * c-c++-common/goacc/reduction-4.c: Likewise. > * c-c++-common/goacc/reduction-8.c: Likewise. > * g++.dg/goacc/firstprivate-mappings-1.C: Likewise. > * g++.dg/gomp/target-lambda-1.C: Likewise. > * g++.dg/gomp/target-this-3.C: Likewise. > * g++.dg/gomp/target-this-4.C: Likewise. > * gfortran.dg/goacc/common-block-3.f90: Likewise. > * gfortran.dg/goacc/loop-tree-1.f90: Likewise. > * gfortran.dg/goacc/private-explicit-kernels-1.f95: Likewise. > * gfortran.dg/goacc/private-predetermined-kernels-1.f95: Likewise. > > libgomp/ChangeLog: > > * target.c (gomp_map_vars_existing): Add 'bool implicit' parameter,= add > implicit map handling to allow a "superset" existing map as valid c= ase. > (get_kind): Adjust to filter out GOMP_MAP_IMPLICIT bits in return v= alue. > (get_implicit): New function to extract implicit status. > (gomp_map_fields_existing): Adjust arguments in calls to > gomp_map_vars_existing, and add uses of get_implicit. > (gomp_map_vars_internal): Likewise. > > * testsuite/libgomp.c-c++-common/target-implicit-map-1.c: New test. > > > > From a70b5b1aa8b3d32f6728dbfcfc00b0cff8c5219d Mon Sep 17 00:00:00 2001 > From: Chung-Lin Tang > Date: Wed, 5 May 2021 08:11:19 -0700 > Subject: [PATCH] OpenMP 5.0: Implement relaxation of implicit map vs. exi= sting > device mappings > > This patch implements relaxing the requirements when a map with the impli= cit > attribute encounters an overlapping existing map. As the OpenMP 5.0 spec > describes on page 320, lines 18-27 (and 5.1 spec, page 352, lines 13-22): > > "If a single contiguous part of the original storage of a list item with = an > implicit data-mapping attribute has corresponding storage in the device = data > environment prior to a task encountering the construct that is associate= d with > the map clause, only that part of the original storage will have corresp= onding > storage in the device data environment as a result of the map clause." > > Also tracked in the OpenMP spec context as issue #1463: > https://github.com/OpenMP/spec/issues/1463 > > 2021-05-05 Chung-Lin Tang > > include/ChangeLog: > > * gomp-constants.h (GOMP_MAP_IMPLICIT): New special map kind bits v= alue. > (GOMP_MAP_FLAG_SPECIAL_BITS): Define helper mask for whole set of > special map kind bits. > (GOMP_MAP_NONCONTIG_ARRAY_P): Adjust test for non-contiguous array = map > kind bits to be more specific. > (GOMP_MAP_IMPLICIT_P): New predicate macro for implicit map kinds. > > gcc/ChangeLog: > > * tree.h (OMP_CLAUSE_MAP_IMPLICIT_P): New access macro for 'implici= t' > bit, using 'base.deprecated_flag' field of tree_node. > * tree-pretty-print.c (dump_omp_clause): Add support for printing > implicit attribute in tree dumping. > * gimplify.c (gimplify_adjust_omp_clauses_1): > Set OMP_CLAUSE_MAP_IMPLICIT_P to 1 if map clause is implicitly crea= ted. > (gimplify_adjust_omp_clauses): Adjust place of adding implicitly cr= eated > clauses, from simple append, to starting of list, after non-map cla= uses. > * omp-low.c (lower_omp_target): Add GOMP_MAP_IMPLICIT bits into kin= d > values passed to libgomp for implicit maps. > > gcc/testsuite/ChangeLog: > > * c-c++-common/gomp/target-implicit-map-1.c: New test. > * c-c++-common/goacc/combined-reduction.c: Adjust scan test pattern= . > * c-c++-common/goacc/firstprivate-mappings-1.c: Likewise. > * c-c++-common/goacc/mdc-1.c: Likewise. > * c-c++-common/goacc/reduction-1.c: Likewise. > * c-c++-common/goacc/reduction-2.c: Likewise. > * c-c++-common/goacc/reduction-3.c: Likewise. > * c-c++-common/goacc/reduction-4.c: Likewise. > * c-c++-common/goacc/reduction-8.c: Likewise. > * g++.dg/goacc/firstprivate-mappings-1.C: Likewise. > * g++.dg/gomp/target-lambda-1.C: Likewise. > * g++.dg/gomp/target-this-3.C: Likewise. > * g++.dg/gomp/target-this-4.C: Likewise. > * gfortran.dg/goacc/common-block-3.f90: Likewise. > * gfortran.dg/goacc/loop-tree-1.f90: Likewise. > * gfortran.dg/goacc/private-explicit-kernels-1.f95: Likewise. > * gfortran.dg/goacc/private-predetermined-kernels-1.f95: Likewise. > > libgomp/ChangeLog: > > * target.c (gomp_map_vars_existing): Add 'bool implicit' parameter,= add > implicit map handling to allow a "superset" existing map as valid c= ase. > (get_kind): Adjust to filter out GOMP_MAP_IMPLICIT bits in return v= alue. > (get_implicit): New function to extract implicit status. > (gomp_map_fields_existing): Adjust arguments in calls to > gomp_map_vars_existing, and add uses of get_implicit. > (gomp_map_vars_internal): Likewise. > > * testsuite/libgomp.c-c++-common/target-implicit-map-1.c: New test. > --- > gcc/gimplify.c | 11 ++- > gcc/omp-low.c | 13 ++++ > .../c-c++-common/goacc/combined-reduction.c | 4 +- > .../c-c++-common/goacc/firstprivate-mappings-1.c | 6 +- > gcc/testsuite/c-c++-common/goacc/mdc-1.c | 2 +- > gcc/testsuite/c-c++-common/goacc/reduction-1.c | 4 +- > gcc/testsuite/c-c++-common/goacc/reduction-2.c | 4 +- > gcc/testsuite/c-c++-common/goacc/reduction-3.c | 4 +- > gcc/testsuite/c-c++-common/goacc/reduction-4.c | 4 +- > gcc/testsuite/c-c++-common/goacc/reduction-8.c | 10 +-- > .../c-c++-common/gomp/target-implicit-map-1.c | 39 +++++++++++ > .../g++.dg/goacc/firstprivate-mappings-1.C | 2 +- > gcc/testsuite/g++.dg/gomp/target-lambda-1.C | 6 +- > gcc/testsuite/g++.dg/gomp/target-this-3.C | 4 +- > gcc/testsuite/g++.dg/gomp/target-this-4.C | 4 +- > gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 | 8 +-- > gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90 | 2 +- > .../goacc/private-explicit-kernels-1.f95 | 4 +- > .../goacc/private-predetermined-kernels-1.f95 | 4 +- > gcc/tree-pretty-print.c | 3 + > gcc/tree.h | 4 ++ > include/gomp-constants.h | 17 ++++- > libgomp/target.c | 78 ++++++++++++++++= ------ > .../libgomp.c-c++-common/target-implicit-map-1.c | 31 +++++++++ > 24 files changed, 211 insertions(+), 57 deletions(-) > create mode 100644 gcc/testsuite/c-c++-common/gomp/target-implicit-map-1= .c > create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-implici= t-map-1.c > > diff --git a/gcc/gimplify.c b/gcc/gimplify.c > index 91aa15d..ba071e8 100644 > --- a/gcc/gimplify.c > +++ b/gcc/gimplify.c > @@ -10579,6 +10579,7 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n,= void *data) > gcc_unreachable (); > } > OMP_CLAUSE_SET_MAP_KIND (clause, kind); > + OMP_CLAUSE_MAP_IMPLICIT_P (clause) =3D 1; > if (DECL_SIZE (decl) > && TREE_CODE (DECL_SIZE (decl)) !=3D INTEGER_CST) > { > @@ -11158,9 +11159,15 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, = gimple_seq body, tree *list_p, > list_p =3D &OMP_CLAUSE_CHAIN (c); > } > > - /* Add in any implicit data sharing. */ > + /* Add in any implicit data sharing. Implicit clauses are added at the= start > + of the clause list, but after any non-map clauses. */ > struct gimplify_adjust_omp_clauses_data data; > - data.list_p =3D list_p; > + tree *implicit_add_list_p =3D orig_list_p; > + while (*implicit_add_list_p > + && OMP_CLAUSE_CODE (*implicit_add_list_p) !=3D OMP_CLAUSE_MAP) > + implicit_add_list_p =3D &OMP_CLAUSE_CHAIN (*implicit_add_list_p); > + > + data.list_p =3D implicit_add_list_p; > data.pre_p =3D pre_p; > splay_tree_foreach (ctx->variables, gimplify_adjust_omp_clauses_1, &da= ta); > > diff --git a/gcc/omp-low.c b/gcc/omp-low.c > index 66519ad..64b7c19 100644 > --- a/gcc/omp-low.c > +++ b/gcc/omp-low.c > @@ -12920,6 +12920,19 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, o= mp_context *ctx) > else if (integer_nonzerop (s)) > tkind_zero =3D tkind; > } > + if (tkind_zero =3D=3D tkind > + && OMP_CLAUSE_MAP_IMPLICIT_P (c) > + && (((tkind & GOMP_MAP_FLAG_SPECIAL_BITS) > + & ~GOMP_MAP_IMPLICIT) > + =3D=3D 0)) > + { > + /* If this is an implicit map, and the GOMP_MAP_IMPLICI= T > + bits are not interfered by other special bit encodin= gs, > + then turn the GOMP_IMPLICIT_BIT flag on for the runt= ime > + to see. */ > + tkind |=3D GOMP_MAP_IMPLICIT; > + tkind_zero =3D tkind; > + } > break; > case OMP_CLAUSE_FIRSTPRIVATE: > gcc_checking_assert (is_gimple_omp_oacc (ctx->stmt)); > diff --git a/gcc/testsuite/c-c++-common/goacc/combined-reduction.c b/gcc/= testsuite/c-c++-common/goacc/combined-reduction.c > index ecf23f5..fa67e08 100644 > --- a/gcc/testsuite/c-c++-common/goacc/combined-reduction.c > +++ b/gcc/testsuite/c-c++-common/goacc/combined-reduction.c > @@ -23,7 +23,7 @@ main () > return 0; > } > > -/* { dg-final { scan-tree-dump-times "omp target oacc_parallel reduction= .+:v1. map.tofrom:v1" 1 "gimple" } } */ > +/* { dg-final { scan-tree-dump-times "omp target oacc_parallel reduction= .+:v1. firstprivate.n. map.tofrom:v1" 1 "gimple" } } */ > /* { dg-final { scan-tree-dump-times "acc loop reduction.+:v1. private.i= ." 1 "gimple" } } */ > -/* { dg-final { scan-tree-dump-times "omp target oacc_kernels map.force_= tofrom:n .len: 4.. map.force_tofrom:v1 .len: 4.." 1 "gimple" } } */ > +/* { dg-final { scan-tree-dump-times "omp target oacc_kernels map.force_= tofrom:n .len: 4..implicit.. map.force_tofrom:v1 .len: 4..implicit.." 1 "gi= mple" } } */ > /* { dg-final { scan-tree-dump-times "acc loop reduction.+:v1. private.i= ." 1 "gimple" } } */ > diff --git a/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c b= /gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c > index 7987bea..f43e4b4 100644 > --- a/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c > +++ b/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c > @@ -419,12 +419,12 @@ vla (int array_li) > copyout (array_so) > /* The gimplifier has created an implicit 'firstprivate' clause for th= e array > length. > - { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel m= ap\(from:array_so \[len: 4\]\) firstprivate\(array_li.[0-9]+\)} omplower { = target { ! c++ } } } } > - { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel m= ap\(from:array_so \[len: 4\]\) firstprivate\(} omplower { target { c++ } } = } } > + { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel f= irstprivate\(array_li.[0-9]+\) map\(from:array_so \[len: 4\]\)} omplower { = target { ! c++ } } } } > + { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel f= irstprivate\([^)]+\) map\(from:array_so \[len: 4\]\)} omplower { target { c= ++ } } } } > (C++ computes an intermediate value, so can't scan for 'firstprivat= e(array_li)'.) */ > /* For C, non-LP64, the gimplifier has also created a mapping for the = array > itself; PR90859. > - { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel m= ap\(from:array_so \[len: 4\]\) firstprivate\(array_li.[0-9]+\) map\(tofrom:= \(\*array.[0-9]+\) \[len: D\.[0-9]+\]\) map\(firstprivate:array \[pointer a= ssign, bias: 0\]\) \[} omplower { target { c && { ! lp64 } } } } } */ > + { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel f= irstprivate\(array_li.[0-9]+\) map\(tofrom:\(\*array.[0-9]+\) \[len: D\.[0-= 9]+\]\[implicit\]\) map\(firstprivate:array \[pointer assign, bias: 0\]\) m= ap\(from:array_so \[len: 4\]\) \[} omplower { target { c && { ! lp64 } } } = } } */ > { > array_so =3D sizeof array; > } > diff --git a/gcc/testsuite/c-c++-common/goacc/mdc-1.c b/gcc/testsuite/c-c= ++-common/goacc/mdc-1.c > index 337c1f7..9f43de4 100644 > --- a/gcc/testsuite/c-c++-common/goacc/mdc-1.c > +++ b/gcc/testsuite/c-c++-common/goacc/mdc-1.c > @@ -45,7 +45,7 @@ t1 () > > /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_= data map.to:s .len: 32.." 1 "omplower" } } */ > /* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.to= from:.z .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len: 8.. map.tofro= m:._1 .len: 40.. map.attach:s.a .bias: 0.." 1 "omplower" } } */ > -/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel ma= p.attach:s.e .bias: 0.. map.tofrom:s .len: 32" 1 "omplower" } } */ > +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel ma= p.tofrom:s .len: 32..implicit.. map.attach:s.e .bias: 0.." 1 "omplower" } }= */ > /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_= data map.attach:a .bias: 0.." 1 "omplower" } } */ > /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_= data map.detach:a .bias: 0.." 1 "omplower" } } */ > /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_= data map.to:a .len: 8.." 1 "omplower" } } */ > diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-1.c b/gcc/testsui= te/c-c++-common/goacc/reduction-1.c > index 35bfc86..d9e3c38 100644 > --- a/gcc/testsuite/c-c++-common/goacc/reduction-1.c > +++ b/gcc/testsuite/c-c++-common/goacc/reduction-1.c > @@ -68,5 +68,5 @@ main(void) > } > > /* Check that default copy maps are generated for loop reductions. */ > -/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9= \]+\\\]\\)" 7 "gimple" } } */ > -/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-= 9\]+\\\]\\)" 2 "gimple" } } */ > +/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9= \]+\\\]\\\[implicit\\\]\\)" 7 "gimple" } } */ > +/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-= 9\]+\\\]\\\[implicit\\\]\\)" 2 "gimple" } } */ > diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-2.c b/gcc/testsui= te/c-c++-common/goacc/reduction-2.c > index 9dba035..18dc03c 100644 > --- a/gcc/testsuite/c-c++-common/goacc/reduction-2.c > +++ b/gcc/testsuite/c-c++-common/goacc/reduction-2.c > @@ -50,5 +50,5 @@ main(void) > } > > /* Check that default copy maps are generated for loop reductions. */ > -/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9= \]+\\\]\\)" 4 "gimple" } } */ > -/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-= 9\]+\\\]\\)" 2 "gimple" } } */ > +/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9= \]+\\\]\\\[implicit\\\]\\)" 4 "gimple" } } */ > +/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-= 9\]+\\\]\\\[implicit\\\]\\)" 2 "gimple" } } */ > diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-3.c b/gcc/testsui= te/c-c++-common/goacc/reduction-3.c > index 669cd43..2311d4b 100644 > --- a/gcc/testsuite/c-c++-common/goacc/reduction-3.c > +++ b/gcc/testsuite/c-c++-common/goacc/reduction-3.c > @@ -50,5 +50,5 @@ main(void) > } > > /* Check that default copy maps are generated for loop reductions. */ > -/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9= \]+\\\]\\)" 4 "gimple" } } */ > -/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-= 9\]+\\\]\\)" 2 "gimple" } } */ > +/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9= \]+\\\]\\\[implicit\\\]\\)" 4 "gimple" } } */ > +/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-= 9\]+\\\]\\\[implicit\\\]\\)" 2 "gimple" } } */ > diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-4.c b/gcc/testsui= te/c-c++-common/goacc/reduction-4.c > index 5c3dfb1..57823f8 100644 > --- a/gcc/testsuite/c-c++-common/goacc/reduction-4.c > +++ b/gcc/testsuite/c-c++-common/goacc/reduction-4.c > @@ -38,5 +38,5 @@ main(void) > } > > /* Check that default copy maps are generated for loop reductions. */ > -/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9= \]+\\\]\\)" 2 "gimple" } } */ > -/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-= 9\]+\\\]\\)" 2 "gimple" } } */ > +/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9= \]+\\\]\\\[implicit\\\]\\)" 2 "gimple" } } */ > +/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-= 9\]+\\\]\\\[implicit\\\]\\)" 2 "gimple" } } */ > diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-8.c b/gcc/testsui= te/c-c++-common/goacc/reduction-8.c > index 8a0283f..8494e59 100644 > --- a/gcc/testsuite/c-c++-common/goacc/reduction-8.c > +++ b/gcc/testsuite/c-c++-common/goacc/reduction-8.c > @@ -87,8 +87,10 @@ main(void) > > /* Check that default copy maps are generated for loop reductions. */ > /* { dg-final { scan-tree-dump-times "reduction..:result. map.tofrom:res= ult .len: 4.." 1 "gimple" } } */ > -/* { dg-final { scan-tree-dump-times "oacc_parallel map.tofrom:result .l= en: 4.." 2 "gimple" } } */ > -/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000.. firs= tprivate.result." 3 "gimple" } } */ > -/* { dg-final { scan-tree-dump-times "map.tofrom:result .len: 4.. map.to= from:array .len: 4000.." 1 "gimple" } } */ > -/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000.. map.= force_tofrom:result .len: 4.." 1 "gimple" } } */ > + > +/* { dg-final { scan-tree-dump-times "oacc_parallel map.tofrom:result .l= en: 4..implicit.." 1 "gimple" } } */ > + > +/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000..impli= cit.. firstprivate.result." 3 "gimple" } } */ > +/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000..impli= cit.. map.tofrom:result .len: 4.." 1 "gimple" } } */ > +/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000..impli= cit.. map.force_tofrom:result .len: 4..implicit.." 1 "gimple" } } */ > > diff --git a/gcc/testsuite/c-c++-common/gomp/target-implicit-map-1.c b/gc= c/testsuite/c-c++-common/gomp/target-implicit-map-1.c > new file mode 100644 > index 0000000..52944fd > --- /dev/null > +++ b/gcc/testsuite/c-c++-common/gomp/target-implicit-map-1.c > @@ -0,0 +1,39 @@ > +/* { dg-do compile } */ > +/* { dg-additional-options "-fdump-tree-gimple" } */ > +#ifdef __cplusplus > +extern "C" > +#else > +extern > +#endif > +void abort (void); > + > +int > +main (void) > +{ > + #define N 5 > + int array[N][N]; > + > + for (int i =3D 0; i < N; i++) > + { > + #pragma omp target enter data map(alloc: array[i:1][0:N]) > + > + #pragma omp target > + for (int j =3D 0; j < N; j++) > + array[i][j] =3D i * 10 + j; > + > + #pragma omp target exit data map(from: array[i:1][0:N]) > + } > + > + for (int i =3D 0; i < N; i++) > + for (int j =3D 0; j < N; j++) > + if (array[i][j] !=3D i + j) > + abort (); > + > + return 0; > +} > + > +/* { dg-final { scan-tree-dump {#pragma omp target enter data map\(alloc= :array\[[^]]+\]\[0\] \[len: [0-9]+\]\)} "gimple" } } */ > + > +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstpriv= ate\(i\) map\(tofrom:array \[len: [0-9]+\]\[implicit\]\)} "gimple" } } */ > + > +/* { dg-final { scan-tree-dump {#pragma omp target exit data map\(from:a= rray\[[^]]+\]\[0\] \[len: [0-9]+\]\)} "gimple" } } */ > diff --git a/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C b/gcc/t= estsuite/g++.dg/goacc/firstprivate-mappings-1.C > index 1b1badb..99a3bd4 100644 > --- a/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C > +++ b/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C > @@ -416,7 +416,7 @@ vla (int &array_li) > copyout (array_so) > /* The gimplifier has created an implicit 'firstprivate' clause for th= e array > length. > - { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel m= ap\(from:array_so \[len: 4\]\) firstprivate\(} omplower } } > + { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel f= irstprivate\([^)]+\) map\(from:array_so \[len: 4\]\)} omplower } } > (C++ computes an intermediate value, so can't scan for 'firstprivat= e(array_li)'.) */ > { > array_so =3D sizeof array; > diff --git a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C b/gcc/testsuite/= g++.dg/gomp/target-lambda-1.C > index 7dceef8..e5a24d7 100644 > --- a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C > +++ b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C > @@ -87,8 +87,8 @@ int main (void) > return 0; > } > > -/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:\= *__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, = bias: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[bias:= 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[= len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer= :__closure->__this \[pointer assign, bias: 0\]\) map\(attach_zero_length_ar= ray_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) m= ap\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstpri= vate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer as= sign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[po= inter assign, bias: 0\]\)} "gimple" } } */ > +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstpriv= ate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer ass= ign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[poi= nter assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstpr= ivate:__closure \[pointer assign, bias: 0\]\) map\(attach_zero_length_array= _section:__closure->__iptr \[bias: 0\]\) map\(struct:\*__closure \[len: 1\]= \) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[l= en: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias= : 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) ma= p\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(all= oc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */ > > -/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:l= oop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data1 \= [bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(end\) firstpri= vate\(begin\)} "gimple" } } */ > +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstpriv= ate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(attach= _zero_length_array_section:loop\.__data1 \[bias: 0\]\) map\(alloc:\*_[0-9]+= \[len: 0\]\)} "gimple" } } */ > > -/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:l= oop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data2 \= [bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(end\) firstpri= vate\(begin\)} "gimple" } } */ > +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstpriv= ate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(attach= _zero_length_array_section:loop\.__data2 \[bias: 0\]\) map\(alloc:\*_[0-9]+= \[len: 0\]\)} "gimple" } } */ > diff --git a/gcc/testsuite/g++.dg/gomp/target-this-3.C b/gcc/testsuite/g+= +.dg/gomp/target-this-3.C > index 08568f9..2755b4b 100644 > --- a/gcc/testsuite/g++.dg/gomp/target-this-3.C > +++ b/gcc/testsuite/g++.dg/gomp/target-this-3.C > @@ -100,6 +100,6 @@ int main (void) > return 0; > } > > -/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofr= om:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: = 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bia= s: 0\]\) map\(attach:this->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-= 9]+\]\) map\(alloc:\*_[0-9+] \[len: 0\]\) firstprivate\(n\)} "gimple" } } *= / > +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstpriv= ate\(n\) map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[poin= ter assign, bias: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length = array section, bias: 0\]\) map\(attach:this->refptr \[bias: 0\]\) map\(from= :mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9+] \[len: 0\]\)} "gimple" } } *= / > > -/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofr= om:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: = 0\]\) map\(attach_zero_length_array_section:this->ptr \[bias: 0\]\) map\(fr= om:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\= (n\)} "gimple" } } */ > +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstpriv= ate\(n\) map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[poin= ter assign, bias: 0\]\) map\(attach_zero_length_array_section:this->ptr \[b= ias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0= \]\)} "gimple" } } */ > diff --git a/gcc/testsuite/g++.dg/gomp/target-this-4.C b/gcc/testsuite/g+= +.dg/gomp/target-this-4.C > index 3b2d581..3703762 100644 > --- a/gcc/testsuite/g++.dg/gomp/target-this-4.C > +++ b/gcc/testsuite/g++.dg/gomp/target-this-4.C > @@ -102,6 +102,6 @@ int main (void) > return 0; > } > > -/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:\= *__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, = bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__th= is \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_po= inter:__closure->__this \[pointer assign, bias: 0\]\) map\(attach_zero_leng= th_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: 1\]\) m= ap\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(n\) map\(alloc:MEM.* \[len: = 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */ > +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstpriv= ate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer ass= ign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__= closure \[pointer assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) = map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len:= [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0= \]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(= from:mapped \[len: 1\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */ > > -/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:\= *__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, = bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__th= is \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_po= inter:__closure->__this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+ = \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:_[0-9]= +->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0= -9]+ \[len: 0\]\) firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firs= tprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */ > +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstpriv= ate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer ass= ign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__= closure \[pointer assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) = map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len:= [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0= \]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias= : 0\]\) map\(attach:_[0-9]+->refptr \[bias: 0\]\) map\(from:mapped \[len: [= 0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */ > diff --git a/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 b/gcc/tes= tsuite/gfortran.dg/goacc/common-block-3.f90 > index e43d376..e9f169f 100644 > --- a/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 > +++ b/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 > @@ -33,10 +33,10 @@ end program main > > ! { dg-final { scan-tree-dump-times "omp target oacc_data_kernels .*map\= \(tofrom:x \\\[len: 400\\\]\\)" 1 "omplower" } } > ! { dg-final { scan-tree-dump-times "omp target oacc_data_kernels .*map\= \(tofrom:y \\\[len: 400\\\]\\\)" 1 "omplower" } } > -! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(for= ce_present:x \\\[len: 400\\\]\\)" 1 "omplower" } } > -! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(for= ce_present:y \\\[len: 400\\\]\\\)" 1 "omplower" } } > -! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(for= ce_tofrom:i \\\[len: 4\\\]\\)" 1 "omplower" } } > -! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(for= ce_tofrom:c \\\[len: 4\\\]\\)" 1 "omplower" } } > +! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(for= ce_present:x \\\[len: 400\\\]\\\[implicit\\\]\\)" 1 "omplower" } } > +! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(for= ce_present:y \\\[len: 400\\\]\\\[implicit\\\]\\\)" 1 "omplower" } } > +! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(for= ce_tofrom:i \\\[len: 4\\\]\\\[implicit\\\]\\)" 1 "omplower" } } > +! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(for= ce_tofrom:c \\\[len: 4\\\]\\\[implicit\\\]\\)" 1 "omplower" } } > > ! Expecting no mapping of un-referenced common-blocks variables > > diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90 b/gcc/testsu= ite/gfortran.dg/goacc/loop-tree-1.f90 > index 150f930..4cdfc55 100644 > --- a/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90 > +++ b/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90 > @@ -44,4 +44,4 @@ end program test > > ! { dg-final { scan-tree-dump-times "private\\(m\\)" 1 "original" } } > ! { dg-final { scan-tree-dump-times "reduction\\(\\+:sum\\)" 1 "original= " } } > -! { dg-final { scan-tree-dump-times "map\\(tofrom:sum \\\[len: \[0-9\]+\= \\]\\)" 1 "gimple" } } > +! { dg-final { scan-tree-dump-times "map\\(tofrom:sum \\\[len: \[0-9\]+\= \\]\\\[implicit\\\]\\)" 1 "gimple" } } > diff --git a/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f= 95 b/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95 > index 0c47045..fef5126 100644 > --- a/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95 > +++ b/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95 > @@ -83,7 +83,7 @@ program test > !$acc kernels ! Explicit "private(i2_2_s)" clause cannot be specified = here. > ! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "original"= { xfail *-*-* } } } ! PR90067 > ! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "gimple" {= xfail *-*-* } } } ! PR90067 > - ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels m= ap\\(force_tofrom:i2_2_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-*= } } } ! PR90067 > + ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels m= ap\\(force_tofrom:i2_2_s \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 0 "gimpl= e" { xfail *-*-* } } } ! PR90067 > do i2_2_s =3D 1, 100 > !$acc loop private(j2_2_s) independent > ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_= 2_s\\) independent" 1 "original" } } > @@ -234,7 +234,7 @@ program test > !$acc kernels ! Explicit "private(i3_5_s)" clause cannot be specified = here. > ! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "original"= { xfail *-*-* } } } ! PR90067 > ! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "gimple" {= xfail *-*-* } } } ! PR90067 > - ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels m= ap\\(force_tofrom:i3_5_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-*= } } } ! PR90067 > + ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels m= ap\\(force_tofrom:i3_5_s \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 0 "gimpl= e" { xfail *-*-* } } } ! PR90067 > do i3_5_s =3D 1, 100 > !$acc loop private(j3_5_s) independent > ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_= 5_s\\) independent" 1 "original" } } > diff --git a/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernel= s-1.f95 b/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f= 95 > index 3357a20..38459cf 100644 > --- a/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95 > +++ b/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95 > @@ -83,7 +83,7 @@ program test > !$acc kernels > ! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "original"= { xfail *-*-* } } } ! PR90067 > ! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "gimple" {= xfail *-*-* } } } ! PR90067 > - ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels m= ap\\(force_tofrom:i2_2_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-*= } } } ! PR90067 > + ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels m= ap\\(force_tofrom:i2_2_s \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 0 "gimpl= e" { xfail *-*-* } } } ! PR90067 > do i2_2_s =3D 1, 100 > !$acc loop independent > ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_= 2_s\\) independent" 1 "original" } } > @@ -234,7 +234,7 @@ program test > !$acc kernels > ! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "original"= { xfail *-*-* } } } ! PR90067 > ! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "gimple" {= xfail *-*-* } } } ! PR90067 > - ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels m= ap\\(force_tofrom:i3_5_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-*= } } } ! PR90067 > + ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels m= ap\\(force_tofrom:i3_5_s \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 0 "gimpl= e" { xfail *-*-* } } } ! PR90067 > do i3_5_s =3D 1, 100 > !$acc loop independent > ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_= 5_s\\) independent" 1 "original" } } > diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c > index 261cc9d..4cd4606 100644 > --- a/gcc/tree-pretty-print.c > +++ b/gcc/tree-pretty-print.c > @@ -946,6 +946,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int= spc, dump_flags_t flags) > spc, flags, false); > pp_right_bracket (pp); > } > + if (OMP_CLAUSE_CODE (clause) =3D=3D OMP_CLAUSE_MAP > + && OMP_CLAUSE_MAP_IMPLICIT_P (clause)) > + pp_string (pp, "[implicit]"); > pp_right_paren (pp); > break; > > diff --git a/gcc/tree.h b/gcc/tree.h > index 8d9829c..647c5ba 100644 > --- a/gcc/tree.h > +++ b/gcc/tree.h > @@ -1637,6 +1637,10 @@ class auto_suppress_location_wrappers > variable. */ > #define OMP_CLAUSE_MAP_IN_REDUCTION(NODE) \ > TREE_PRIVATE (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)) > +/* Nonzero if this map clause was created through implicit data-mapping > + rules. */ > +#define OMP_CLAUSE_MAP_IMPLICIT_P(NODE) \ > + (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.deprecated_flag= ) > > /* True on an OMP_CLAUSE_USE_DEVICE_PTR with an OpenACC 'if_present' > clause. */ > diff --git a/include/gomp-constants.h b/include/gomp-constants.h > index b8efb30..33cfcb9 100644 > --- a/include/gomp-constants.h > +++ b/include/gomp-constants.h > @@ -46,6 +46,16 @@ > | GOMP_MAP_FLAG_SPECIAL_0) > #define GOMP_MAP_DEEP_COPY (GOMP_MAP_FLAG_SPECIAL_4 \ > | GOMP_MAP_FLAG_SPECIAL_2) > +/* This value indicates the map was created implicitly according to > + OpenMP rules. */ > +#define GOMP_MAP_IMPLICIT (GOMP_MAP_FLAG_SPECIAL_3 \ > + | GOMP_MAP_FLAG_SPECIAL_4) > +/* Mask for entire set of special map kind bits. */ > +#define GOMP_MAP_FLAG_SPECIAL_BITS (GOMP_MAP_FLAG_SPECIAL_0 \ > + | GOMP_MAP_FLAG_SPECIAL_1 \ > + | GOMP_MAP_FLAG_SPECIAL_2 \ > + | GOMP_MAP_FLAG_SPECIAL_3 \ > + | GOMP_MAP_FLAG_SPECIAL_4) > /* Flag to force a specific behavior (or else, trigger a run-time error)= . */ > #define GOMP_MAP_FLAG_FORCE (1 << 7) > > @@ -225,7 +235,12 @@ enum gomp_map_kind > (GOMP_MAP_ALWAYS_TO_P (X) || ((X) =3D=3D GOMP_MAP_ALWAYS_FROM)) > > #define GOMP_MAP_NONCONTIG_ARRAY_P(X) \ > - ((X) & GOMP_MAP_NONCONTIG_ARRAY) > + (((X) & GOMP_MAP_FLAG_SPECIAL_BITS) =3D=3D GOMP_MAP_NONCONTIG_ARRAY = \ > + || (X) =3D=3D GOMP_MAP_NONCONTIG_ARRAY_FORCE_PRESENT) > + > +#define GOMP_MAP_IMPLICIT_P(X) \ > + (((X) & GOMP_MAP_FLAG_SPECIAL_BITS) =3D=3D GOMP_MAP_IMPLICIT) > + > > /* Asynchronous behavior. Keep in sync with > libgomp/{openacc.h,openacc.f90,openacc_lib.h}:acc_async_t. */ > diff --git a/libgomp/target.c b/libgomp/target.c > index 9c75826..ecda2ef 100644 > --- a/libgomp/target.c > +++ b/libgomp/target.c > @@ -510,7 +510,7 @@ static inline void > gomp_map_vars_existing (struct gomp_device_descr *devicep, > struct goacc_asyncqueue *aq, splay_tree_key oldn, > splay_tree_key newn, struct target_var_desc *tgt_va= r, > - unsigned char kind, bool always_to_flag, > + unsigned char kind, bool always_to_flag, bool impli= cit, > struct gomp_coalesce_buf *cbuf, > htab_t *refcount_set) > { > @@ -522,11 +522,22 @@ gomp_map_vars_existing (struct gomp_device_descr *d= evicep, > tgt_var->always_copy_from =3D GOMP_MAP_ALWAYS_FROM_P (kind); > tgt_var->is_attach =3D false; > tgt_var->offset =3D newn->host_start - oldn->host_start; > - tgt_var->length =3D newn->host_end - newn->host_start; > + > + /* For implicit maps, old contained in new is valid. */ > + bool implicit_subset =3D (implicit > + && newn->host_start <=3D oldn->host_start > + && oldn->host_end <=3D newn->host_end); > + if (implicit_subset) > + tgt_var->length =3D oldn->host_end - oldn->host_start; > + else > + tgt_var->length =3D newn->host_end - newn->host_start; > > if ((kind & GOMP_MAP_FLAG_FORCE) > - || oldn->host_start > newn->host_start > - || oldn->host_end < newn->host_end) > + /* For implicit maps, old contained in new is valid. */ > + || !(implicit_subset > + /* Otherwise, new contained inside old is considered valid. */ > + || (oldn->host_start <=3D newn->host_start > + && newn->host_end <=3D oldn->host_end))) > { > gomp_mutex_unlock (&devicep->lock); > gomp_fatal ("Trying to map into device [%p..%p) object when " > @@ -536,11 +547,17 @@ gomp_map_vars_existing (struct gomp_device_descr *d= evicep, > } > > if (GOMP_MAP_ALWAYS_TO_P (kind) || always_to_flag) > - gomp_copy_host2dev (devicep, aq, > - (void *) (oldn->tgt->tgt_start + oldn->tgt_offset > - + newn->host_start - oldn->host_start), > - (void *) newn->host_start, > - newn->host_end - newn->host_start, false, cbuf); > + { > + /* Implicit + always should not happen. If this does occur, below > + address/length adjustment is a TODO. */ > + assert (!implicit_subset); > + > + gomp_copy_host2dev (devicep, aq, > + (void *) (oldn->tgt->tgt_start + oldn->tgt_offset > + + newn->host_start - oldn->host_start), > + (void *) newn->host_start, > + newn->host_end - newn->host_start, false, cbuf); > + } > > gomp_increment_refcount (oldn, refcount_set); > } > @@ -548,8 +565,24 @@ gomp_map_vars_existing (struct gomp_device_descr *de= vicep, > static int > get_kind (bool short_mapkind, void *kinds, int idx) > { > - return short_mapkind ? ((unsigned short *) kinds)[idx] > - : ((unsigned char *) kinds)[idx]; > + int val =3D (short_mapkind > + ? ((unsigned short *) kinds)[idx] > + : ((unsigned char *) kinds)[idx]); > + > + if (GOMP_MAP_IMPLICIT_P (val)) > + val &=3D ~GOMP_MAP_IMPLICIT; > + return val; > +} > + > + > +static bool > +get_implicit (bool short_mapkind, void *kinds, int idx) > +{ > + int val =3D (short_mapkind > + ? ((unsigned short *) kinds)[idx] > + : ((unsigned char *) kinds)[idx]); > + > + return GOMP_MAP_IMPLICIT_P (val); > } > > static void > @@ -612,6 +645,7 @@ gomp_map_fields_existing (struct target_mem_desc *tgt= , > struct splay_tree_s *mem_map =3D &devicep->mem_map; > struct splay_tree_key_s cur_node; > int kind; > + bool implicit; > const bool short_mapkind =3D true; > const int typemask =3D short_mapkind ? 0xff : 0x7; > > @@ -619,12 +653,14 @@ gomp_map_fields_existing (struct target_mem_desc *t= gt, > cur_node.host_end =3D cur_node.host_start + sizes[i]; > splay_tree_key n2 =3D splay_tree_lookup (mem_map, &cur_node); > kind =3D get_kind (short_mapkind, kinds, i); > + implicit =3D get_implicit (short_mapkind, kinds, i); > if (n2 > && n2->tgt =3D=3D n->tgt > && n2->host_start - n->host_start =3D=3D n2->tgt_offset - n->tgt_o= ffset) > { > gomp_map_vars_existing (devicep, aq, n2, &cur_node, &tgt->list[i], > - kind & typemask, false, cbuf, refcount_set); > + kind & typemask, false, implicit, cbuf, > + refcount_set); > return; > } > if (sizes[i] =3D=3D 0) > @@ -640,7 +676,8 @@ gomp_map_fields_existing (struct target_mem_desc *tgt= , > =3D=3D n2->tgt_offset - n->tgt_offset) > { > gomp_map_vars_existing (devicep, aq, n2, &cur_node, &tgt->lis= t[i], > - kind & typemask, false, cbuf, refcoun= t_set); > + kind & typemask, false, implicit, cbu= f, > + refcount_set); > return; > } > } > @@ -652,7 +689,8 @@ gomp_map_fields_existing (struct target_mem_desc *tgt= , > && n2->host_start - n->host_start =3D=3D n2->tgt_offset - n->tgt_= offset) > { > gomp_map_vars_existing (devicep, aq, n2, &cur_node, &tgt->list[i]= , > - kind & typemask, false, cbuf, refcount_se= t); > + kind & typemask, false, implicit, cbuf, > + refcount_set); > return; > } > } > @@ -898,6 +936,7 @@ gomp_map_vars_internal (struct gomp_device_descr *dev= icep, > for (i =3D 0; i < mapnum; i++) > { > int kind =3D get_kind (short_mapkind, kinds, i); > + bool implicit =3D get_implicit (short_mapkind, kinds, i); > if (hostaddrs[i] =3D=3D NULL > || (kind & typemask) =3D=3D GOMP_MAP_FIRSTPRIVATE_INT) > { > @@ -1104,8 +1143,8 @@ gomp_map_vars_internal (struct gomp_device_descr *d= evicep, > } > } > gomp_map_vars_existing (devicep, aq, n, &cur_node, &tgt->list[i], > - kind & typemask, always_to_cnt > 0, NULL, > - refcount_set); > + kind & typemask, always_to_cnt > 0, impli= cit, > + NULL, refcount_set); > i +=3D always_to_cnt; > } > else > @@ -1182,7 +1221,7 @@ gomp_map_vars_internal (struct gomp_device_descr *d= evicep, > { > assert (n->refcount !=3D REFCOUNT_LINK); > gomp_map_vars_existing (devicep, aq, n, &cur_node, row_de= sc, > - kind & typemask, false, > + kind & typemask, false, false, > /* TODO: cbuf? */ NULL, refcount_= set); > } > else > @@ -1312,6 +1351,7 @@ gomp_map_vars_internal (struct gomp_device_descr *d= evicep, > else if (tgt->list[i].key =3D=3D NULL) > { > int kind =3D get_kind (short_mapkind, kinds, i); > + bool implicit =3D get_implicit (short_mapkind, kinds, i); > if (hostaddrs[i] =3D=3D NULL) > continue; > switch (kind & typemask) > @@ -1483,7 +1523,7 @@ gomp_map_vars_internal (struct gomp_device_descr *d= evicep, > splay_tree_key n =3D splay_tree_lookup (mem_map, k); > if (n && n->refcount !=3D REFCOUNT_LINK) > gomp_map_vars_existing (devicep, aq, n, k, &tgt->list[i], > - kind & typemask, false, cbufp, > + kind & typemask, false, implicit, cbu= fp, > refcount_set); > else > { > @@ -1702,7 +1742,7 @@ gomp_map_vars_internal (struct gomp_device_descr *d= evicep, > { > assert (k->refcount !=3D REFCOUNT_LINK); > gomp_map_vars_existing (devicep, aq, k, &cur_node, ro= w_desc, > - kind & typemask, false, > + kind & typemask, false, false= , > cbufp, refcount_set); > } > else > diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-1= .c b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-1.c > new file mode 100644 > index 0000000..f2e7293 > --- /dev/null > +++ b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-1.c > @@ -0,0 +1,31 @@ > +#ifdef __cplusplus > +extern "C" > +#else > +extern > +#endif > +void abort (void); > + > +int > +main (void) > +{ > + #define N 5 > + int array[N][N]; > + > + for (int i =3D 0; i < N; i++) > + { > + #pragma omp target enter data map(alloc: array[i:1][0:N]) > + > + #pragma omp target > + for (int j =3D 0; j < N; j++) > + array[i][j] =3D i + j; > + > + #pragma omp target exit data map(from: array[i:1][0:N]) > + } > + > + for (int i =3D 0; i < N; i++) > + for (int j =3D 0; j < N; j++) > + if (array[i][j] !=3D i + j) > + abort (); > + > + return 0; > +} ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 M=C3=BCnchen R= egistergericht M=C3=BCnchen HRB 106955, Gesch=C3=A4ftsf=C3=BChrer: Thomas H= eurung, Frank Th=C3=BCrauf --=-=-= Content-Type: text/x-diff Content-Disposition: inline; filename="0001-Fix-up-c-c-common-goacc-firstprivate-mappings-1.og10.patch" >From c51cc3b96f0b562deaffcfbcc51043aed216801a Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Fri, 7 May 2021 10:57:36 +0200 Subject: [PATCH] Fix up 'c-c++-common/goacc/firstprivate-mappings-1.c' for C, non-LP64 Follow-up to recent og10 commit a70b5b1aa8b3d32f6728dbfcfc00b0cff8c5219d "OpenMP 5.0: Implement relaxation of implicit map vs. existing device mappings". gcc/testsuite/ * c-c++-common/goacc/firstprivate-mappings-1.c: Fix up for C, non-LP64. --- gcc/testsuite/ChangeLog.omp | 5 +++++ gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c | 4 ++-- 2 files changed, 7 insertions(+), 2 deletions(-) diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp index 45277a37fd2..b18615a1603 100644 --- a/gcc/testsuite/ChangeLog.omp +++ b/gcc/testsuite/ChangeLog.omp @@ -1,3 +1,8 @@ +2021-05-07 Thomas Schwinge + + * c-c++-common/goacc/firstprivate-mappings-1.c: Fix up for C, + non-LP64. + 2021-04-11 Hafiz Abid Qadeer Backport from mainline diff --git a/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c b/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c index f43e4b46cb6..ab09dee3d37 100644 --- a/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c +++ b/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c @@ -419,8 +419,8 @@ vla (int array_li) copyout (array_so) /* The gimplifier has created an implicit 'firstprivate' clause for the array length. - { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel firstprivate\(array_li.[0-9]+\) map\(from:array_so \[len: 4\]\)} omplower { target { ! c++ } } } } - { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel firstprivate\([^)]+\) map\(from:array_so \[len: 4\]\)} omplower { target { c++ } } } } + { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel firstprivate\(array_li.[0-9]+\) map\(from:array_so \[len: 4\]\) \[} omplower { target { c && lp64 } } } } + { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel firstprivate\(D\.[0-9]+\) map\(from:array_so \[len: 4\]\) \[} omplower { target { c++ } } } } (C++ computes an intermediate value, so can't scan for 'firstprivate(array_li)'.) */ /* For C, non-LP64, the gimplifier has also created a mapping for the array itself; PR90859. -- 2.30.2 --=-=-=--