From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa4.mentor.iphmx.com (esa4.mentor.iphmx.com [68.232.137.252]) by sourceware.org (Postfix) with ESMTPS id 13ED63857001 for ; Mon, 7 Jun 2021 11:28:44 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 13ED63857001 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com IronPort-SDR: kPeehq/VR2jSZyDtvUjx6rH+sHIpfj7/KqpmTZbu9XJ4paDWqnnq+iWM/AUzUur8YToCODwb8Z WMR4a/QRazGz5kRK7yuVc/NsNyPiHnBSidXwPyIZHL/OWiUhf0tYU+KIQeRZS3K3iFlkVpjQ6y EK5mpJBufMrFpX8oK8191l3FZ/k5iiAJLVTAKy18hs7GNeZy5pywtP4rZkal3bb9DKNhdRDl5A kLOFaLxno1zZHOESwP6DZNpIMYJ1dTxI2faSUY15hd8i7zELDj8O/cgWOzp0G73f3c/OaQI8m5 l/I= X-IronPort-AV: E=Sophos;i="5.83,254,1616486400"; d="scan'208";a="62237530" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa4.mentor.iphmx.com with ESMTP; 07 Jun 2021 03:28:42 -0800 IronPort-SDR: C+yvdk0lUuAOsFWYUZRt8+6QcmczK+RRd+j8GalekRzMzPpitbaz4qW6YoHV72cFuabqI/lQQe aqbeQiPa4+ojQ9ssw74m3ToZAIcjQ+7LKWpxLArxH6l6BAAeb55psg8yoWGa3+AmoDtNfEP3ex w3Z9hsIJP7c1coTa7tiSYbTTv2LX0kf2Ulo5Gaaz1KYKTzFKpWwnkWqHHyYSKQUk6jZqdRwZyt UDquyZBr7OjrKcQ8u6QG9zi0CzbvWk8SM8gRd7WXcKxmg3JUHLVXqtmfxL7gmxfkRo4a+C2jBV HP0= From: Thomas Schwinge To: Chung-Lin Tang CC: , Jakub Jelinek , "Catherine Moore" , Tobias Burnus Subject: Re: [PATCH, OpenMP 5.0] Implement relaxation of implicit map vs. existing device mappings (for mainline trunk) In-Reply-To: References: <87tuneu3f4.fsf@euler.schwinge.homeip.net> User-Agent: Notmuch/0.29.3+94~g74c3f1b (https://notmuchmail.org) Emacs/27.1 (x86_64-pc-linux-gnu) Date: Mon, 7 Jun 2021 13:28:33 +0200 Message-ID: <87sg1tsyla.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 Content-Type: text/plain; charset="utf-8" Content-Transfer-Encoding: quoted-printable X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) To svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) X-Spam-Status: No, score=-12.4 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_SHORT, 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: Mon, 07 Jun 2021 11:28:55 -0000 Hi Chung-Lin! On 2021-05-14T21:20:25+0800, Chung-Lin Tang wrote= : > This is a version of patch https://gcc.gnu.org/pipermail/gcc-patches/2021= -May/569665.html > for mainline trunk. Related to the discussion in that thread, , please keep this disabled for OpenACC, for the time being. I do like the general idea (but haven't reviewed in detail the implementation), but this needs some more thought (and additional changes) for OpenACC, also related to other patches that are to be upstreamed. Does your 'OMP_CLAUSE_MAP_IMPLICIT_P': /* 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_fla= g) ... need to be integrated/refactored regarding the 'OMP_CLAUSE_MAP_IMPLICIT' that Jakub recently added in commit r12-1109-gc94424b0ed786ec92b6904da69af8b5243b34fdc "openmp: Fix up handling of reduction clause on constructs combined with target [PR99928]": /* Nonzero on map clauses added implicitly for reduction clauses on com= bined or composite constructs. They shall be removed if there is an expli= cit map clause. */ #define OMP_CLAUSE_MAP_IMPLICIT(NODE) \ (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.default_def_fl= ag) Gr=C3=BC=C3=9Fe Thomas > 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 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. > > There are also some additions to print the implicit attribute during tree= pretty-printing, for that > reason some scan tests were updated. > > 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". > > 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. > > 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 on x86_64-linux with nvptx offloading. > Was already pushed to devel/omp/gcc-10 a while ago, asking for approval f= or mainline trunk. > > Chung-Lin > > 2021-05-14 Chung-Lin Tang > > include/ChangeLog: > > * gomp-constants.h (GOMP_MAP_FLAG_SPECIAL_3): Define special bit ma= cro. > (GOMP_MAP_IMPLICIT): New special map kind bits value. > (GOMP_MAP_FLAG_SPECIAL_BITS): Define helper mask for whole set of > special map kind bits. > (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. > diff --git a/gcc/gimplify.c b/gcc/gimplify.c > index e790f08b23f..69c4a8e0a0a 100644 > --- a/gcc/gimplify.c > +++ b/gcc/gimplify.c > @@ -10374,6 +10374,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) > { > @@ -10971,9 +10972,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 cadca7e201f..e8fdd2741bb 100644 > --- a/gcc/omp-low.c > +++ b/gcc/omp-low.c > @@ -12498,6 +12498,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 ecf23f59d66..fa67e085c86 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 7987beaed9a..5134ef6ed6c 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,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\(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++ } } = } } > - (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\(from:array_so \[len: 4\]\) \[} omplower= } } */ > { > 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 337c1f7cc77..9f43de4f776 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 35bfc868708..d9e3c380b8e 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 9dba035adb6..18dc03c93ac 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 669cd438113..2311d4b0adb 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 5c3dfb19172..57823f8898f 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/gomp/target-implicit-map-1.c b/gc= c/testsuite/c-c++-common/gomp/target-implicit-map-1.c > new file mode 100644 > index 00000000000..52944fdc65a > --- /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 1b1badb1a90..99a3bd472f7 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/gfortran.dg/goacc/common-block-3.f90 b/gcc/tes= tsuite/gfortran.dg/goacc/common-block-3.f90 > index 5defe2ea85d..dd98afe4fb1 100644 > --- a/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 > +++ b/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 > @@ -30,10 +30,10 @@ end program main > ! { dg-final { scan-tree-dump-times "omp target oacc_parallel .*map\\(to= from:b \\\[len: 400\\\]\\\)" 1 "omplower" } } > ! { dg-final { scan-tree-dump-times "omp target oacc_parallel .*map\\(to= from:c \\\[len: 4\\\]\\)" 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\\(tof= rom:x \\\[len: 400\\\]\\)" 1 "omplower" } } > -! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(tof= rom:y \\\[len: 400\\\]\\\)" 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_tofrom:i \\\[len: 4\\\]\\\[implicit\\\]\\)" 1 "omplower" } } > +! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(tof= rom:x \\\[len: 400\\\]\\\[implicit\\\]\\)" 1 "omplower" } } > +! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(tof= rom:y \\\[len: 400\\\]\\\[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 150f9304e46..4cdfc5556b7 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 5d563d226b0..eedd986c7b9 100644 > --- a/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95 > +++ b/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95 > @@ -82,7 +82,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" } } > @@ -231,7 +231,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 12a7854526a..24bc0e73906 100644 > --- a/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95 > +++ b/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95 > @@ -82,7 +82,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" } } > @@ -231,7 +231,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 0a575eb9dad..8f026d332ea 100644 > --- a/gcc/tree-pretty-print.c > +++ b/gcc/tree-pretty-print.c > @@ -929,6 +929,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 784452ca490..83b920a35ff 100644 > --- a/gcc/tree.h > +++ b/gcc/tree.h > @@ -1644,6 +1644,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 6e163b02560..6e65b6437b3 100644 > --- a/include/gomp-constants.h > +++ b/include/gomp-constants.h > @@ -40,11 +40,22 @@ > #define GOMP_MAP_FLAG_SPECIAL_0 (1 << 2) > #define GOMP_MAP_FLAG_SPECIAL_1 (1 << 3) > #define GOMP_MAP_FLAG_SPECIAL_2 (1 << 4) > +#define GOMP_MAP_FLAG_SPECIAL_3 (1 << 5) > #define GOMP_MAP_FLAG_SPECIAL_4 (1 << 6) > #define GOMP_MAP_FLAG_SPECIAL (GOMP_MAP_FLAG_SPECIAL_1 \ > | 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) > > @@ -186,6 +197,9 @@ enum gomp_map_kind > #define GOMP_MAP_ALWAYS_P(X) \ > (GOMP_MAP_ALWAYS_TO_P (X) || ((X) =3D=3D GOMP_MAP_ALWAYS_FROM)) > > +#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 2150e5d79b2..b836e3d7f80 100644 > --- a/libgomp/target.c > +++ b/libgomp/target.c > @@ -368,7 +368,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) > { > assert (kind !=3D GOMP_MAP_ATTACH); > @@ -378,11 +378,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 " > @@ -392,11 +403,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, 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, cbuf); > + } > > if (oldn->refcount !=3D REFCOUNT_INFINITY) > oldn->refcount++; > @@ -405,8 +422,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 > @@ -459,6 +492,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; > > @@ -466,12 +500,13 @@ 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); > + kind & typemask, false, implicit, cbuf); > return; > } > if (sizes[i] =3D=3D 0) > @@ -487,7 +522,7 @@ 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); > + kind & typemask, false, implicit, cbu= f); > return; > } > } > @@ -499,7 +534,7 @@ 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); > + kind & typemask, false, implicit, cbuf); > return; > } > } > @@ -729,6 +764,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) > { > @@ -909,7 +945,8 @@ gomp_map_vars_internal (struct gomp_device_descr *dev= icep, > } > } > gomp_map_vars_existing (devicep, aq, n, &cur_node, &tgt->list[i], > - kind & typemask, always_to_cnt > 0, NULL)= ; > + kind & typemask, always_to_cnt > 0, impli= cit, > + NULL); > i +=3D always_to_cnt; > } > else > @@ -1078,6 +1115,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) > @@ -1236,7 +1274,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); > else > { > k->aux =3D NULL; > 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 00000000000..f2e72936862 > --- /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