| Message ID | 09ae7f35-41fd-c83e-2054-18ad4433d5e3@gmail.com |
|---|---|
| State | New |
| Headers | show |
| Series | [OG10,OpenMP,5.0,committed] Implement relaxation of implicit map vs. existing device mappings | expand |
Hi Chung-Lin! On 2021-05-05T23:17:25+0800, Chung-Lin Tang via Gcc-patches <gcc-patches@gcc.gnu.org> wrote: > This patch implements relaxing the requirements when a map with the implicit 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 prior 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 implicitly 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 == GOMP_MAP_FLAG_SPECIAL_3|GOMP_MAP_FLAG_SPECIAL_4 > in include/gomp-constants.h for the runtime, but I've intentionally avoided 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 map success if an existing > map is a proper subset of the new map, if the new map is implicit. Straightforward 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 = 0; i < N; i++) > rec.ptr[i] += 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 before 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üße 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* be mapping its base-pointer > (although an attachment attempt should exist), while in current GCC behavior, 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 really exhibit the above mentioned > behavior yet. I have included it as part of this patch because the "[implicit]" tree printing requires > modifying many gimple scan tests already, so including the test modifications together seems more > manageable patch-wise. > > Tested with no regressions, and pushed to devel/omp/gcc-10. Will be submitting a mainline trunk version later. > > Chung-Lin > > 2021-05-05 Chung-Lin Tang <cltang@codesourcery.com> > > include/ChangeLog: > > * gomp-constants.h (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_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 'implicit' > 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 created. > (gimplify_adjust_omp_clauses): Adjust place of adding implicitly created > clauses, from simple append, to starting of list, after non-map clauses. > * omp-low.c (lower_omp_target): Add GOMP_MAP_IMPLICIT bits into kind > 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 case. > (get_kind): Adjust to filter out GOMP_MAP_IMPLICIT bits in return value. > (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 <cltang@codesourcery.com> > Date: Wed, 5 May 2021 08:11:19 -0700 > Subject: [PATCH] OpenMP 5.0: Implement relaxation of implicit map vs. existing > device mappings > > This patch implements relaxing the requirements when a map with the implicit > 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 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 > > 2021-05-05 Chung-Lin Tang <cltang@codesourcery.com> > > include/ChangeLog: > > * gomp-constants.h (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_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 'implicit' > 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 created. > (gimplify_adjust_omp_clauses): Adjust place of adding implicitly created > clauses, from simple append, to starting of list, after non-map clauses. > * omp-low.c (lower_omp_target): Add GOMP_MAP_IMPLICIT bits into kind > 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 case. > (get_kind): Adjust to filter out GOMP_MAP_IMPLICIT bits in return value. > (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-implicit-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) = 1; > if (DECL_SIZE (decl) > && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST) > { > @@ -11158,9 +11159,15 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, > list_p = &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 = list_p; > + tree *implicit_add_list_p = orig_list_p; > + while (*implicit_add_list_p > + && OMP_CLAUSE_CODE (*implicit_add_list_p) != OMP_CLAUSE_MAP) > + implicit_add_list_p = &OMP_CLAUSE_CHAIN (*implicit_add_list_p); > + > + data.list_p = implicit_add_list_p; > data.pre_p = pre_p; > splay_tree_foreach (ctx->variables, gimplify_adjust_omp_clauses_1, &data); > > 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, omp_context *ctx) > else if (integer_nonzerop (s)) > tkind_zero = tkind; > } > + if (tkind_zero == tkind > + && OMP_CLAUSE_MAP_IMPLICIT_P (c) > + && (((tkind & GOMP_MAP_FLAG_SPECIAL_BITS) > + & ~GOMP_MAP_IMPLICIT) > + == 0)) > + { > + /* If this is an implicit map, and the GOMP_MAP_IMPLICIT > + bits are not interfered by other special bit encodings, > + then turn the GOMP_IMPLICIT_BIT flag on for the runtime > + to see. */ > + tkind |= GOMP_MAP_IMPLICIT; > + tkind_zero = 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 "gimple" } } */ > /* { 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 the array > length. > - { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel map\(from:array_so \[len: 4\]\) firstprivate\(array_li.[0-9]+\)} omplower { target { ! c++ } } } } > - { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel map\(from:array_so \[len: 4\]\) firstprivate\(} 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++ } } } } > + { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel firstprivate\([^)]+\) 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. > - { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel map\(from:array_so \[len: 4\]\) firstprivate\(array_li.[0-9]+\) map\(tofrom:\(\*array.[0-9]+\) \[len: D\.[0-9]+\]\) map\(firstprivate:array \[pointer assign, bias: 0\]\) \[} omplower { target { c && { ! lp64 } } } } } */ > + { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel firstprivate\(array_li.[0-9]+\) map\(tofrom:\(\*array.[0-9]+\) \[len: D\.[0-9]+\]\[implicit\]\) map\(firstprivate:array \[pointer assign, bias: 0\]\) map\(from:array_so \[len: 4\]\) \[} omplower { target { c && { ! lp64 } } } } } */ > { > array_so = 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.tofrom:.z .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len: 8.. map.tofrom:._1 .len: 40.. map.attach:s.a .bias: 0.." 1 "omplower" } } */ > -/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.attach:s.e .bias: 0.. map.tofrom:s .len: 32" 1 "omplower" } } */ > +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.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/testsuite/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/testsuite/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/testsuite/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/testsuite/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/testsuite/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:result .len: 4.." 1 "gimple" } } */ > -/* { dg-final { scan-tree-dump-times "oacc_parallel map.tofrom:result .len: 4.." 2 "gimple" } } */ > -/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000.. firstprivate.result." 3 "gimple" } } */ > -/* { dg-final { scan-tree-dump-times "map.tofrom:result .len: 4.. map.tofrom: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 .len: 4..implicit.." 1 "gimple" } } */ > + > +/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000..implicit.. firstprivate.result." 3 "gimple" } } */ > +/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000..implicit.. map.tofrom:result .len: 4.." 1 "gimple" } } */ > +/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000..implicit.. map.force_tofrom:result .len: 4..implicit.." 1 "gimple" } } */ > > diff --git a/gcc/testsuite/c-c++-common/gomp/target-implicit-map-1.c b/gcc/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 = 0; i < N; i++) > + { > + #pragma omp target enter data map(alloc: array[i:1][0:N]) > + > + #pragma omp target > + for (int j = 0; j < N; j++) > + array[i][j] = i * 10 + j; > + > + #pragma omp target exit data map(from: array[i:1][0:N]) > + } > + > + for (int i = 0; i < N; i++) > + for (int j = 0; j < N; j++) > + if (array[i][j] != 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.* firstprivate\(i\) map\(tofrom:array \[len: [0-9]+\]\[implicit\]\)} "gimple" } } */ > + > +/* { dg-final { scan-tree-dump {#pragma omp target exit data map\(from:array\[[^]]+\]\[0\] \[len: [0-9]+\]\)} "gimple" } } */ > diff --git a/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C b/gcc/testsuite/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 the array > length. > - { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel map\(from:array_so \[len: 4\]\) firstprivate\(} omplower } } > + { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel firstprivate\([^)]+\) map\(from:array_so \[len: 4\]\)} omplower } } > (C++ computes an intermediate value, so can't scan for 'firstprivate(array_li)'.) */ > { > array_so = 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_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */ > +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) 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_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */ > > -/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(end\) firstprivate\(begin\)} "gimple" } } */ > +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(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:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data2 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(end\) firstprivate\(begin\)} "gimple" } } */ > +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(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\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer 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\]\) firstprivate\(n\)} "gimple" } } */ > +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer 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\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:this->ptr \[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.* firstprivate\(n\) map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:this->ptr \[bias: 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->__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\]\) 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.* firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, 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->__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\]\) 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.* firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, 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/testsuite/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\\(force_present:x \\\[len: 400\\\]\\)" 1 "omplower" } } > -! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:y \\\[len: 400\\\]\\\)" 1 "omplower" } } > -! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_tofrom:i \\\[len: 4\\\]\\)" 1 "omplower" } } > -! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_tofrom:c \\\[len: 4\\\]\\)" 1 "omplower" } } > +! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:x \\\[len: 400\\\]\\\[implicit\\\]\\)" 1 "omplower" } } > +! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:y \\\[len: 400\\\]\\\[implicit\\\]\\\)" 1 "omplower" } } > +! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_tofrom:i \\\[len: 4\\\]\\\[implicit\\\]\\)" 1 "omplower" } } > +! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_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/testsuite/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.f95 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 map\\(force_tofrom:i2_2_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067 > + ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i2_2_s \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067 > do i2_2_s = 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 map\\(force_tofrom:i3_5_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067 > + ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i3_5_s \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067 > do i3_5_s = 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-kernels-1.f95 b/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95 > 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 map\\(force_tofrom:i2_2_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067 > + ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i2_2_s \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067 > do i2_2_s = 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 map\\(force_tofrom:i3_5_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067 > + ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i3_5_s \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067 > do i3_5_s = 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) == 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) == GOMP_MAP_ALWAYS_FROM)) > > #define GOMP_MAP_NONCONTIG_ARRAY_P(X) \ > - ((X) & GOMP_MAP_NONCONTIG_ARRAY) > + (((X) & GOMP_MAP_FLAG_SPECIAL_BITS) == GOMP_MAP_NONCONTIG_ARRAY \ > + || (X) == GOMP_MAP_NONCONTIG_ARRAY_FORCE_PRESENT) > + > +#define GOMP_MAP_IMPLICIT_P(X) \ > + (((X) & GOMP_MAP_FLAG_SPECIAL_BITS) == 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_var, > - unsigned char kind, bool always_to_flag, > + unsigned char kind, bool always_to_flag, bool implicit, > struct gomp_coalesce_buf *cbuf, > htab_t *refcount_set) > { > @@ -522,11 +522,22 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep, > tgt_var->always_copy_from = GOMP_MAP_ALWAYS_FROM_P (kind); > tgt_var->is_attach = false; > tgt_var->offset = newn->host_start - oldn->host_start; > - tgt_var->length = newn->host_end - newn->host_start; > + > + /* For implicit maps, old contained in new is valid. */ > + bool implicit_subset = (implicit > + && newn->host_start <= oldn->host_start > + && oldn->host_end <= newn->host_end); > + if (implicit_subset) > + tgt_var->length = oldn->host_end - oldn->host_start; > + else > + tgt_var->length = 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 <= newn->host_start > + && newn->host_end <= 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 *devicep, > } > > 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 *devicep, > static int > get_kind (bool short_mapkind, void *kinds, int idx) > { > - return short_mapkind ? ((unsigned short *) kinds)[idx] > - : ((unsigned char *) kinds)[idx]; > + int val = (short_mapkind > + ? ((unsigned short *) kinds)[idx] > + : ((unsigned char *) kinds)[idx]); > + > + if (GOMP_MAP_IMPLICIT_P (val)) > + val &= ~GOMP_MAP_IMPLICIT; > + return val; > +} > + > + > +static bool > +get_implicit (bool short_mapkind, void *kinds, int idx) > +{ > + int val = (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 = &devicep->mem_map; > struct splay_tree_key_s cur_node; > int kind; > + bool implicit; > const bool short_mapkind = true; > const int typemask = short_mapkind ? 0xff : 0x7; > > @@ -619,12 +653,14 @@ gomp_map_fields_existing (struct target_mem_desc *tgt, > cur_node.host_end = cur_node.host_start + sizes[i]; > splay_tree_key n2 = splay_tree_lookup (mem_map, &cur_node); > kind = get_kind (short_mapkind, kinds, i); > + implicit = get_implicit (short_mapkind, kinds, i); > if (n2 > && n2->tgt == n->tgt > && n2->host_start - n->host_start == n2->tgt_offset - n->tgt_offset) > { > 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] == 0) > @@ -640,7 +676,8 @@ gomp_map_fields_existing (struct target_mem_desc *tgt, > == n2->tgt_offset - n->tgt_offset) > { > 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; > } > } > @@ -652,7 +689,8 @@ gomp_map_fields_existing (struct target_mem_desc *tgt, > && n2->host_start - n->host_start == n2->tgt_offset - n->tgt_offset) > { > 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; > } > } > @@ -898,6 +936,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, > for (i = 0; i < mapnum; i++) > { > int kind = get_kind (short_mapkind, kinds, i); > + bool implicit = get_implicit (short_mapkind, kinds, i); > if (hostaddrs[i] == NULL > || (kind & typemask) == GOMP_MAP_FIRSTPRIVATE_INT) > { > @@ -1104,8 +1143,8 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, > } > } > 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, implicit, > + NULL, refcount_set); > i += always_to_cnt; > } > else > @@ -1182,7 +1221,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, > { > assert (n->refcount != REFCOUNT_LINK); > gomp_map_vars_existing (devicep, aq, n, &cur_node, row_desc, > - 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 *devicep, > else if (tgt->list[i].key == NULL) > { > int kind = get_kind (short_mapkind, kinds, i); > + bool implicit = get_implicit (short_mapkind, kinds, i); > if (hostaddrs[i] == NULL) > continue; > switch (kind & typemask) > @@ -1483,7 +1523,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, > splay_tree_key n = splay_tree_lookup (mem_map, k); > if (n && n->refcount != REFCOUNT_LINK) > gomp_map_vars_existing (devicep, aq, n, k, &tgt->list[i], > - kind & typemask, false, cbufp, > + kind & typemask, false, implicit, cbufp, > refcount_set); > else > { > @@ -1702,7 +1742,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, > { > assert (k->refcount != REFCOUNT_LINK); > gomp_map_vars_existing (devicep, aq, k, &cur_node, row_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 = 0; i < N; i++) > + { > + #pragma omp target enter data map(alloc: array[i:1][0:N]) > + > + #pragma omp target > + for (int j = 0; j < N; j++) > + array[i][j] = i + j; > + > + #pragma omp target exit data map(from: array[i:1][0:N]) > + } > + > + for (int i = 0; i < N; i++) > + for (int j = 0; j < N; j++) > + if (array[i][j] != i + j) > + abort (); > + > + return 0; > +} ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf
On 2021/5/7 8:35 PM, Thomas Schwinge wrote: > On 2021-05-05T23:17:25+0800, Chung-Lin Tang via Gcc-patches<gcc-patches@gcc.gnu.org> wrote: >> This patch implements relaxing the requirements when a map with the implicit 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? It appears that I have inadvertently enabled it for OpenACC as well! But everything was tested together, so I assume it works okay for that mode as well. The entire set of implicit-specific actions are enabled by the setting of 'OMP_CLAUSE_MAP_IMPLICIT_P (clause) = 1' inside gimplify.c:gimplify_adjust_omp_clauses_1, so in case you want to disable it for OpenACC again, that's where you need to add the guard condition. >> 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. Thanks, I was relying on eyeballing to know where to fix testcases like this; I did fix another similar case, but missed this one. > >> The reason for this is: when combined with other map clauses, for example: >> >> #pragma omp target map(rec.ptr[:N]) >> for (int i = 0; i < N; i++) >> rec.ptr[i] += 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 before 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?) It probably won't affect user level semantics, although we should look out if this change in convention exposes some other bugs. Chung-Lin
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) = 1; if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST) { @@ -11158,9 +11159,15 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, list_p = &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 = list_p; + tree *implicit_add_list_p = orig_list_p; + while (*implicit_add_list_p + && OMP_CLAUSE_CODE (*implicit_add_list_p) != OMP_CLAUSE_MAP) + implicit_add_list_p = &OMP_CLAUSE_CHAIN (*implicit_add_list_p); + + data.list_p = implicit_add_list_p; data.pre_p = pre_p; splay_tree_foreach (ctx->variables, gimplify_adjust_omp_clauses_1, &data); 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, omp_context *ctx) else if (integer_nonzerop (s)) tkind_zero = tkind; } + if (tkind_zero == tkind + && OMP_CLAUSE_MAP_IMPLICIT_P (c) + && (((tkind & GOMP_MAP_FLAG_SPECIAL_BITS) + & ~GOMP_MAP_IMPLICIT) + == 0)) + { + /* If this is an implicit map, and the GOMP_MAP_IMPLICIT + bits are not interfered by other special bit encodings, + then turn the GOMP_IMPLICIT_BIT flag on for the runtime + to see. */ + tkind |= GOMP_MAP_IMPLICIT; + tkind_zero = 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 "gimple" } } */ /* { 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 the array length. - { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel map\(from:array_so \[len: 4\]\) firstprivate\(array_li.[0-9]+\)} omplower { target { ! c++ } } } } - { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel map\(from:array_so \[len: 4\]\) firstprivate\(} 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++ } } } } + { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel firstprivate\([^)]+\) 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. - { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel map\(from:array_so \[len: 4\]\) firstprivate\(array_li.[0-9]+\) map\(tofrom:\(\*array.[0-9]+\) \[len: D\.[0-9]+\]\) map\(firstprivate:array \[pointer assign, bias: 0\]\) \[} omplower { target { c && { ! lp64 } } } } } */ + { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel firstprivate\(array_li.[0-9]+\) map\(tofrom:\(\*array.[0-9]+\) \[len: D\.[0-9]+\]\[implicit\]\) map\(firstprivate:array \[pointer assign, bias: 0\]\) map\(from:array_so \[len: 4\]\) \[} omplower { target { c && { ! lp64 } } } } } */ { array_so = 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.tofrom:.z .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len: 8.. map.tofrom:._1 .len: 40.. map.attach:s.a .bias: 0.." 1 "omplower" } } */ -/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.attach:s.e .bias: 0.. map.tofrom:s .len: 32" 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.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/testsuite/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/testsuite/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/testsuite/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/testsuite/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/testsuite/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:result .len: 4.." 1 "gimple" } } */ -/* { dg-final { scan-tree-dump-times "oacc_parallel map.tofrom:result .len: 4.." 2 "gimple" } } */ -/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000.. firstprivate.result." 3 "gimple" } } */ -/* { dg-final { scan-tree-dump-times "map.tofrom:result .len: 4.. map.tofrom: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 .len: 4..implicit.." 1 "gimple" } } */ + +/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000..implicit.. firstprivate.result." 3 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000..implicit.. map.tofrom:result .len: 4.." 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000..implicit.. map.force_tofrom:result .len: 4..implicit.." 1 "gimple" } } */ diff --git a/gcc/testsuite/c-c++-common/gomp/target-implicit-map-1.c b/gcc/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 = 0; i < N; i++) + { + #pragma omp target enter data map(alloc: array[i:1][0:N]) + + #pragma omp target + for (int j = 0; j < N; j++) + array[i][j] = i * 10 + j; + + #pragma omp target exit data map(from: array[i:1][0:N]) + } + + for (int i = 0; i < N; i++) + for (int j = 0; j < N; j++) + if (array[i][j] != 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.* firstprivate\(i\) map\(tofrom:array \[len: [0-9]+\]\[implicit\]\)} "gimple" } } */ + +/* { dg-final { scan-tree-dump {#pragma omp target exit data map\(from:array\[[^]]+\]\[0\] \[len: [0-9]+\]\)} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C b/gcc/testsuite/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 the array length. - { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel map\(from:array_so \[len: 4\]\) firstprivate\(} omplower } } + { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel firstprivate\([^)]+\) map\(from:array_so \[len: 4\]\)} omplower } } (C++ computes an intermediate value, so can't scan for 'firstprivate(array_li)'.) */ { array_so = 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_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */ +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) 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_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */ -/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(end\) firstprivate\(begin\)} "gimple" } } */ +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(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:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data2 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(end\) firstprivate\(begin\)} "gimple" } } */ +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(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\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer 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\]\) firstprivate\(n\)} "gimple" } } */ +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer 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\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:this->ptr \[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.* firstprivate\(n\) map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:this->ptr \[bias: 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->__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\]\) 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.* firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, 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->__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\]\) 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.* firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, 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/testsuite/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\\(force_present:x \\\[len: 400\\\]\\)" 1 "omplower" } } -! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:y \\\[len: 400\\\]\\\)" 1 "omplower" } } -! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_tofrom:i \\\[len: 4\\\]\\)" 1 "omplower" } } -! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_tofrom:c \\\[len: 4\\\]\\)" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:x \\\[len: 400\\\]\\\[implicit\\\]\\)" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:y \\\[len: 400\\\]\\\[implicit\\\]\\\)" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_tofrom:i \\\[len: 4\\\]\\\[implicit\\\]\\)" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_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/testsuite/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.f95 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 map\\(force_tofrom:i2_2_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067 + ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i2_2_s \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067 do i2_2_s = 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 map\\(force_tofrom:i3_5_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067 + ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i3_5_s \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067 do i3_5_s = 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-kernels-1.f95 b/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95 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 map\\(force_tofrom:i2_2_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067 + ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i2_2_s \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067 do i2_2_s = 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 map\\(force_tofrom:i3_5_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067 + ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i3_5_s \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067 do i3_5_s = 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) == 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) == GOMP_MAP_ALWAYS_FROM)) #define GOMP_MAP_NONCONTIG_ARRAY_P(X) \ - ((X) & GOMP_MAP_NONCONTIG_ARRAY) + (((X) & GOMP_MAP_FLAG_SPECIAL_BITS) == GOMP_MAP_NONCONTIG_ARRAY \ + || (X) == GOMP_MAP_NONCONTIG_ARRAY_FORCE_PRESENT) + +#define GOMP_MAP_IMPLICIT_P(X) \ + (((X) & GOMP_MAP_FLAG_SPECIAL_BITS) == 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_var, - unsigned char kind, bool always_to_flag, + unsigned char kind, bool always_to_flag, bool implicit, struct gomp_coalesce_buf *cbuf, htab_t *refcount_set) { @@ -522,11 +522,22 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep, tgt_var->always_copy_from = GOMP_MAP_ALWAYS_FROM_P (kind); tgt_var->is_attach = false; tgt_var->offset = newn->host_start - oldn->host_start; - tgt_var->length = newn->host_end - newn->host_start; + + /* For implicit maps, old contained in new is valid. */ + bool implicit_subset = (implicit + && newn->host_start <= oldn->host_start + && oldn->host_end <= newn->host_end); + if (implicit_subset) + tgt_var->length = oldn->host_end - oldn->host_start; + else + tgt_var->length = 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 <= newn->host_start + && newn->host_end <= 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 *devicep, } 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 *devicep, static int get_kind (bool short_mapkind, void *kinds, int idx) { - return short_mapkind ? ((unsigned short *) kinds)[idx] - : ((unsigned char *) kinds)[idx]; + int val = (short_mapkind + ? ((unsigned short *) kinds)[idx] + : ((unsigned char *) kinds)[idx]); + + if (GOMP_MAP_IMPLICIT_P (val)) + val &= ~GOMP_MAP_IMPLICIT; + return val; +} + + +static bool +get_implicit (bool short_mapkind, void *kinds, int idx) +{ + int val = (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 = &devicep->mem_map; struct splay_tree_key_s cur_node; int kind; + bool implicit; const bool short_mapkind = true; const int typemask = short_mapkind ? 0xff : 0x7; @@ -619,12 +653,14 @@ gomp_map_fields_existing (struct target_mem_desc *tgt, cur_node.host_end = cur_node.host_start + sizes[i]; splay_tree_key n2 = splay_tree_lookup (mem_map, &cur_node); kind = get_kind (short_mapkind, kinds, i); + implicit = get_implicit (short_mapkind, kinds, i); if (n2 && n2->tgt == n->tgt && n2->host_start - n->host_start == n2->tgt_offset - n->tgt_offset) { 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] == 0) @@ -640,7 +676,8 @@ gomp_map_fields_existing (struct target_mem_desc *tgt, == n2->tgt_offset - n->tgt_offset) { 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; } } @@ -652,7 +689,8 @@ gomp_map_fields_existing (struct target_mem_desc *tgt, && n2->host_start - n->host_start == n2->tgt_offset - n->tgt_offset) { 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; } } @@ -898,6 +936,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, for (i = 0; i < mapnum; i++) { int kind = get_kind (short_mapkind, kinds, i); + bool implicit = get_implicit (short_mapkind, kinds, i); if (hostaddrs[i] == NULL || (kind & typemask) == GOMP_MAP_FIRSTPRIVATE_INT) { @@ -1104,8 +1143,8 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, } } 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, implicit, + NULL, refcount_set); i += always_to_cnt; } else @@ -1182,7 +1221,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, { assert (n->refcount != REFCOUNT_LINK); gomp_map_vars_existing (devicep, aq, n, &cur_node, row_desc, - 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 *devicep, else if (tgt->list[i].key == NULL) { int kind = get_kind (short_mapkind, kinds, i); + bool implicit = get_implicit (short_mapkind, kinds, i); if (hostaddrs[i] == NULL) continue; switch (kind & typemask) @@ -1483,7 +1523,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, splay_tree_key n = splay_tree_lookup (mem_map, k); if (n && n->refcount != REFCOUNT_LINK) gomp_map_vars_existing (devicep, aq, n, k, &tgt->list[i], - kind & typemask, false, cbufp, + kind & typemask, false, implicit, cbufp, refcount_set); else { @@ -1702,7 +1742,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, { assert (k->refcount != REFCOUNT_LINK); gomp_map_vars_existing (devicep, aq, k, &cur_node, row_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 = 0; i < N; i++) + { + #pragma omp target enter data map(alloc: array[i:1][0:N]) + + #pragma omp target + for (int j = 0; j < N; j++) + array[i][j] = i + j; + + #pragma omp target exit data map(from: array[i:1][0:N]) + } + + for (int i = 0; i < N; i++) + for (int j = 0; j < N; j++) + if (array[i][j] != i + j) + abort (); + + return 0; +}
This patch implements relaxing the requirements when a map with the implicit 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 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 implicitly 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 == GOMP_MAP_FLAG_SPECIAL_3|GOMP_MAP_FLAG_SPECIAL_4 in include/gomp-constants.h for the runtime, but I've intentionally avoided 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 map success if an existing map is a proper subset of the new map, if the new map is implicit. Straightforward 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 = 0; i < N; i++) rec.ptr[i] += 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 before 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* be mapping its base-pointer (although an attachment attempt should exist), while in current GCC behavior, 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 really exhibit the above mentioned behavior yet. I have included it as part of this patch because the "[implicit]" tree printing requires modifying many gimple scan tests already, so including the test modifications together seems more manageable patch-wise. Tested with no regressions, and pushed to devel/omp/gcc-10. Will be submitting a mainline trunk version later. Chung-Lin 2021-05-05 Chung-Lin Tang <cltang@codesourcery.com> include/ChangeLog: * gomp-constants.h (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_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 'implicit' 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 created. (gimplify_adjust_omp_clauses): Adjust place of adding implicitly created clauses, from simple append, to starting of list, after non-map clauses. * omp-low.c (lower_omp_target): Add GOMP_MAP_IMPLICIT bits into kind 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 case. (get_kind): Adjust to filter out GOMP_MAP_IMPLICIT bits in return value. (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 <cltang@codesourcery.com> Date: Wed, 5 May 2021 08:11:19 -0700 Subject: [PATCH] OpenMP 5.0: Implement relaxation of implicit map vs. existing device mappings This patch implements relaxing the requirements when a map with the implicit 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 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 2021-05-05 Chung-Lin Tang <cltang@codesourcery.com> include/ChangeLog: * gomp-constants.h (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_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 'implicit' 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 created. (gimplify_adjust_omp_clauses): Adjust place of adding implicitly created clauses, from simple append, to starting of list, after non-map clauses. * omp-low.c (lower_omp_target): Add GOMP_MAP_IMPLICIT bits into kind 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 case. (get_kind): Adjust to filter out GOMP_MAP_IMPLICIT bits in return value. (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-implicit-map-1.c