Hi Chung-Lin! On 2021-05-14T21:20:25+0800, Chung-Lin Tang <clt...@codesourcery.com> wrote: > This is a version of patch > https://gcc.gnu.org/pipermail/gcc-patches/2021-May/569665.html > for mainline trunk.
Related to the discussion in that thread, <87tuneu3f4.fsf@euler.schwinge.homeip.net">http://mid.mail-archive.com/87tuneu3f4.fsf@euler.schwinge.homeip.net>, please keep this disabled for OpenACC, for the time being. I do like the general idea (but haven't reviewed in detail the implementation), but this needs some more thought (and additional changes) for OpenACC, also related to other patches that are to be upstreamed. Does your 'OMP_CLAUSE_MAP_IMPLICIT_P': /* Nonzero if this map clause was created through implicit data-mapping rules. */ #define OMP_CLAUSE_MAP_IMPLICIT_P(NODE) \ (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.deprecated_flag) ... need to be integrated/refactored regarding the 'OMP_CLAUSE_MAP_IMPLICIT' that Jakub recently added in commit r12-1109-gc94424b0ed786ec92b6904da69af8b5243b34fdc "openmp: Fix up handling of reduction clause on constructs combined with target [PR99928]": /* Nonzero on map clauses added implicitly for reduction clauses on combined or composite constructs. They shall be removed if there is an explicit map clause. */ #define OMP_CLAUSE_MAP_IMPLICIT(NODE) \ (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.default_def_flag) Grüße Thomas > 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 on x86_64-linux with nvptx offloading. > Was already pushed to devel/omp/gcc-10 a while ago, asking for approval for > mainline trunk. > > Chung-Lin > > 2021-05-14 Chung-Lin Tang <clt...@codesourcery.com> > > include/ChangeLog: > > * gomp-constants.h (GOMP_MAP_FLAG_SPECIAL_3): Define special bit macro. > (GOMP_MAP_IMPLICIT): New special map kind bits value. > (GOMP_MAP_FLAG_SPECIAL_BITS): Define helper mask for whole set of > special map kind bits. > (GOMP_MAP_IMPLICIT_P): New predicate macro for implicit map kinds. > > gcc/ChangeLog: > > * tree.h (OMP_CLAUSE_MAP_IMPLICIT_P): New access macro for '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. > diff --git a/gcc/gimplify.c b/gcc/gimplify.c > index e790f08b23f..69c4a8e0a0a 100644 > --- a/gcc/gimplify.c > +++ b/gcc/gimplify.c > @@ -10374,6 +10374,7 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, > void *data) > gcc_unreachable (); > } > OMP_CLAUSE_SET_MAP_KIND (clause, kind); > + OMP_CLAUSE_MAP_IMPLICIT_P (clause) = 1; > if (DECL_SIZE (decl) > && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST) > { > @@ -10971,9 +10972,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 cadca7e201f..e8fdd2741bb 100644 > --- a/gcc/omp-low.c > +++ b/gcc/omp-low.c > @@ -12498,6 +12498,19 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, > 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 ecf23f59d66..fa67e085c86 100644 > --- a/gcc/testsuite/c-c++-common/goacc/combined-reduction.c > +++ b/gcc/testsuite/c-c++-common/goacc/combined-reduction.c > @@ -23,7 +23,7 @@ main () > return 0; > } > > -/* { dg-final { scan-tree-dump-times "omp target oacc_parallel > reduction.+:v1. map.tofrom:v1" 1 "gimple" } } */ > +/* { dg-final { scan-tree-dump-times "omp target oacc_parallel > reduction.+:v1. firstprivate.n. map.tofrom:v1" 1 "gimple" } } */ > /* { dg-final { scan-tree-dump-times "acc loop reduction.+:v1. private.i." 1 > "gimple" } } */ > -/* { dg-final { scan-tree-dump-times "omp target oacc_kernels > map.force_tofrom:n .len: 4.. map.force_tofrom:v1 .len: 4.." 1 "gimple" } } */ > +/* { dg-final { scan-tree-dump-times "omp target oacc_kernels > map.force_tofrom:n .len: 4..implicit.. map.force_tofrom:v1 .len: > 4..implicit.." 1 "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 7987beaed9a..5134ef6ed6c 100644 > --- a/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c > +++ b/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c > @@ -419,12 +419,7 @@ vla (int array_li) > copyout (array_so) > /* The gimplifier has created an implicit 'firstprivate' clause for 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++ } } } > } > - (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\(from:array_so \[len: 4\]\) \[} omplower > } } */ > { > 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 337c1f7cc77..9f43de4f776 100644 > --- a/gcc/testsuite/c-c++-common/goacc/mdc-1.c > +++ b/gcc/testsuite/c-c++-common/goacc/mdc-1.c > @@ -45,7 +45,7 @@ t1 () > > /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data > map.to:s .len: 32.." 1 "omplower" } } */ > /* { dg-final { scan-tree-dump-times "pragma omp target oacc_data > map.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 35bfc868708..d9e3c380b8e 100644 > --- a/gcc/testsuite/c-c++-common/goacc/reduction-1.c > +++ b/gcc/testsuite/c-c++-common/goacc/reduction-1.c > @@ -68,5 +68,5 @@ main(void) > } > > /* Check that default copy maps are generated for loop reductions. */ > -/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: > \[0-9\]+\\\]\\)" 7 "gimple" } } */ > -/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: > \[0-9\]+\\\]\\)" 2 "gimple" } } */ > +/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: > \[0-9\]+\\\]\\\[implicit\\\]\\)" 7 "gimple" } } */ > +/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: > \[0-9\]+\\\]\\\[implicit\\\]\\)" 2 "gimple" } } */ > diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-2.c > b/gcc/testsuite/c-c++-common/goacc/reduction-2.c > index 9dba035adb6..18dc03c93ac 100644 > --- a/gcc/testsuite/c-c++-common/goacc/reduction-2.c > +++ b/gcc/testsuite/c-c++-common/goacc/reduction-2.c > @@ -50,5 +50,5 @@ main(void) > } > > /* Check that default copy maps are generated for loop reductions. */ > -/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: > \[0-9\]+\\\]\\)" 4 "gimple" } } */ > -/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: > \[0-9\]+\\\]\\)" 2 "gimple" } } */ > +/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: > \[0-9\]+\\\]\\\[implicit\\\]\\)" 4 "gimple" } } */ > +/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: > \[0-9\]+\\\]\\\[implicit\\\]\\)" 2 "gimple" } } */ > diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-3.c > b/gcc/testsuite/c-c++-common/goacc/reduction-3.c > index 669cd438113..2311d4b0adb 100644 > --- a/gcc/testsuite/c-c++-common/goacc/reduction-3.c > +++ b/gcc/testsuite/c-c++-common/goacc/reduction-3.c > @@ -50,5 +50,5 @@ main(void) > } > > /* Check that default copy maps are generated for loop reductions. */ > -/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: > \[0-9\]+\\\]\\)" 4 "gimple" } } */ > -/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: > \[0-9\]+\\\]\\)" 2 "gimple" } } */ > +/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: > \[0-9\]+\\\]\\\[implicit\\\]\\)" 4 "gimple" } } */ > +/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: > \[0-9\]+\\\]\\\[implicit\\\]\\)" 2 "gimple" } } */ > diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-4.c > b/gcc/testsuite/c-c++-common/goacc/reduction-4.c > index 5c3dfb19172..57823f8898f 100644 > --- a/gcc/testsuite/c-c++-common/goacc/reduction-4.c > +++ b/gcc/testsuite/c-c++-common/goacc/reduction-4.c > @@ -38,5 +38,5 @@ main(void) > } > > /* Check that default copy maps are generated for loop reductions. */ > -/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: > \[0-9\]+\\\]\\)" 2 "gimple" } } */ > -/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: > \[0-9\]+\\\]\\)" 2 "gimple" } } */ > +/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: > \[0-9\]+\\\]\\\[implicit\\\]\\)" 2 "gimple" } } */ > +/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: > \[0-9\]+\\\]\\\[implicit\\\]\\)" 2 "gimple" } } */ > diff --git a/gcc/testsuite/c-c++-common/gomp/target-implicit-map-1.c > b/gcc/testsuite/c-c++-common/gomp/target-implicit-map-1.c > new file mode 100644 > index 00000000000..52944fdc65a > --- /dev/null > +++ b/gcc/testsuite/c-c++-common/gomp/target-implicit-map-1.c > @@ -0,0 +1,39 @@ > +/* { dg-do compile } */ > +/* { dg-additional-options "-fdump-tree-gimple" } */ > +#ifdef __cplusplus > +extern "C" > +#else > +extern > +#endif > +void abort (void); > + > +int > +main (void) > +{ > + #define N 5 > + int array[N][N]; > + > + for (int i = 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 1b1badb1a90..99a3bd472f7 100644 > --- a/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C > +++ b/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C > @@ -416,7 +416,7 @@ vla (int &array_li) > copyout (array_so) > /* The gimplifier has created an implicit 'firstprivate' clause for 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/gfortran.dg/goacc/common-block-3.f90 > b/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 > index 5defe2ea85d..dd98afe4fb1 100644 > --- a/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 > +++ b/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 > @@ -30,10 +30,10 @@ end program main > ! { dg-final { scan-tree-dump-times "omp target oacc_parallel > .*map\\(tofrom:b \\\[len: 400\\\]\\\)" 1 "omplower" } } > ! { dg-final { scan-tree-dump-times "omp target oacc_parallel > .*map\\(tofrom:c \\\[len: 4\\\]\\)" 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\\(tofrom:x \\\[len: 400\\\]\\)" 1 "omplower" } } > -! { dg-final { scan-tree-dump-times "omp target oacc_kernels > .*map\\(tofrom:y \\\[len: 400\\\]\\\)" 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_tofrom:i \\\[len: 4\\\]\\\[implicit\\\]\\)" 1 "omplower" } } > +! { dg-final { scan-tree-dump-times "omp target oacc_kernels > .*map\\(tofrom:x \\\[len: 400\\\]\\\[implicit\\\]\\)" 1 "omplower" } } > +! { dg-final { scan-tree-dump-times "omp target oacc_kernels > .*map\\(tofrom:y \\\[len: 400\\\]\\\[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 150f9304e46..4cdfc5556b7 100644 > --- a/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90 > +++ b/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90 > @@ -44,4 +44,4 @@ end program test > > ! { dg-final { scan-tree-dump-times "private\\(m\\)" 1 "original" } } > ! { dg-final { scan-tree-dump-times "reduction\\(\\+:sum\\)" 1 "original" } } > -! { dg-final { scan-tree-dump-times "map\\(tofrom:sum \\\[len: > \[0-9\]+\\\]\\)" 1 "gimple" } } > +! { dg-final { scan-tree-dump-times "map\\(tofrom:sum \\\[len: > \[0-9\]+\\\]\\\[implicit\\\]\\)" 1 "gimple" } } > diff --git a/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95 > b/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95 > index 5d563d226b0..eedd986c7b9 100644 > --- a/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95 > +++ b/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95 > @@ -82,7 +82,7 @@ program test > !$acc kernels ! Explicit "private(i2_2_s)" clause cannot be specified here. > ! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "original" { > xfail *-*-* } } } ! PR90067 > ! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "gimple" { > xfail *-*-* } } } ! PR90067 > - ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels > 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" } } > @@ -231,7 +231,7 @@ program test > !$acc kernels ! Explicit "private(i3_5_s)" clause cannot be specified here. > ! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "original" { > xfail *-*-* } } } ! PR90067 > ! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "gimple" { > xfail *-*-* } } } ! PR90067 > - ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels > 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 12a7854526a..24bc0e73906 100644 > --- a/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95 > +++ b/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95 > @@ -82,7 +82,7 @@ program test > !$acc kernels > ! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "original" { > xfail *-*-* } } } ! PR90067 > ! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "gimple" { > xfail *-*-* } } } ! PR90067 > - ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels > 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" } } > @@ -231,7 +231,7 @@ program test > !$acc kernels > ! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "original" { > xfail *-*-* } } } ! PR90067 > ! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "gimple" { > xfail *-*-* } } } ! PR90067 > - ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels > 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 0a575eb9dad..8f026d332ea 100644 > --- a/gcc/tree-pretty-print.c > +++ b/gcc/tree-pretty-print.c > @@ -929,6 +929,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int > spc, dump_flags_t flags) > spc, flags, false); > pp_right_bracket (pp); > } > + if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP > + && OMP_CLAUSE_MAP_IMPLICIT_P (clause)) > + pp_string (pp, "[implicit]"); > pp_right_paren (pp); > break; > > diff --git a/gcc/tree.h b/gcc/tree.h > index 784452ca490..83b920a35ff 100644 > --- a/gcc/tree.h > +++ b/gcc/tree.h > @@ -1644,6 +1644,10 @@ class auto_suppress_location_wrappers > variable. */ > #define OMP_CLAUSE_MAP_IN_REDUCTION(NODE) \ > TREE_PRIVATE (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)) > +/* Nonzero if this map clause was created through implicit data-mapping > + rules. */ > +#define OMP_CLAUSE_MAP_IMPLICIT_P(NODE) \ > + (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.deprecated_flag) > > /* True on an OMP_CLAUSE_USE_DEVICE_PTR with an OpenACC 'if_present' > clause. */ > diff --git a/include/gomp-constants.h b/include/gomp-constants.h > index 6e163b02560..6e65b6437b3 100644 > --- a/include/gomp-constants.h > +++ b/include/gomp-constants.h > @@ -40,11 +40,22 @@ > #define GOMP_MAP_FLAG_SPECIAL_0 (1 << 2) > #define GOMP_MAP_FLAG_SPECIAL_1 (1 << 3) > #define GOMP_MAP_FLAG_SPECIAL_2 (1 << 4) > +#define GOMP_MAP_FLAG_SPECIAL_3 (1 << 5) > #define GOMP_MAP_FLAG_SPECIAL_4 (1 << 6) > #define GOMP_MAP_FLAG_SPECIAL (GOMP_MAP_FLAG_SPECIAL_1 \ > | GOMP_MAP_FLAG_SPECIAL_0) > #define GOMP_MAP_DEEP_COPY (GOMP_MAP_FLAG_SPECIAL_4 \ > | GOMP_MAP_FLAG_SPECIAL_2) > +/* This value indicates the map was created implicitly according to > + OpenMP rules. */ > +#define GOMP_MAP_IMPLICIT (GOMP_MAP_FLAG_SPECIAL_3 \ > + | GOMP_MAP_FLAG_SPECIAL_4) > +/* Mask for entire set of special map kind bits. */ > +#define GOMP_MAP_FLAG_SPECIAL_BITS (GOMP_MAP_FLAG_SPECIAL_0 \ > + | GOMP_MAP_FLAG_SPECIAL_1 \ > + | GOMP_MAP_FLAG_SPECIAL_2 \ > + | GOMP_MAP_FLAG_SPECIAL_3 \ > + | GOMP_MAP_FLAG_SPECIAL_4) > /* Flag to force a specific behavior (or else, trigger a run-time error). */ > #define GOMP_MAP_FLAG_FORCE (1 << 7) > > @@ -186,6 +197,9 @@ enum gomp_map_kind > #define GOMP_MAP_ALWAYS_P(X) \ > (GOMP_MAP_ALWAYS_TO_P (X) || ((X) == GOMP_MAP_ALWAYS_FROM)) > > +#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 2150e5d79b2..b836e3d7f80 100644 > --- a/libgomp/target.c > +++ b/libgomp/target.c > @@ -368,7 +368,7 @@ static inline void > gomp_map_vars_existing (struct gomp_device_descr *devicep, > struct goacc_asyncqueue *aq, splay_tree_key oldn, > splay_tree_key newn, struct target_var_desc *tgt_var, > - unsigned char kind, bool always_to_flag, > + unsigned char kind, bool always_to_flag, bool implicit, > struct gomp_coalesce_buf *cbuf) > { > assert (kind != GOMP_MAP_ATTACH); > @@ -378,11 +378,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 " > @@ -392,11 +403,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, cbuf); > + { > + /* Implicit + always should not happen. If this does occur, below > + address/length adjustment is a TODO. */ > + assert (!implicit_subset); > + > + gomp_copy_host2dev (devicep, aq, > + (void *) (oldn->tgt->tgt_start + oldn->tgt_offset > + + newn->host_start - oldn->host_start), > + (void *) newn->host_start, > + newn->host_end - newn->host_start, cbuf); > + } > > if (oldn->refcount != REFCOUNT_INFINITY) > oldn->refcount++; > @@ -405,8 +422,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 > @@ -459,6 +492,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; > > @@ -466,12 +500,13 @@ 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); > + kind & typemask, false, implicit, cbuf); > return; > } > if (sizes[i] == 0) > @@ -487,7 +522,7 @@ 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); > + kind & typemask, false, implicit, cbuf); > return; > } > } > @@ -499,7 +534,7 @@ 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); > + kind & typemask, false, implicit, cbuf); > return; > } > } > @@ -729,6 +764,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) > { > @@ -909,7 +945,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); > + kind & typemask, always_to_cnt > 0, implicit, > + NULL); > i += always_to_cnt; > } > else > @@ -1078,6 +1115,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) > @@ -1236,7 +1274,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); > else > { > k->aux = NULL; > diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-1.c > b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-1.c > new file mode 100644 > index 00000000000..f2e72936862 > --- /dev/null > +++ b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-1.c > @@ -0,0 +1,31 @@ > +#ifdef __cplusplus > +extern "C" > +#else > +extern > +#endif > +void abort (void); > + > +int > +main (void) > +{ > + #define N 5 > + int array[N][N]; > + > + for (int i = 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
