| Message ID | ab750663-5dd0-f6e1-3595-fc3770945ecd@codesourcery.com |
|---|---|
| State | New |
| Headers | show |
| Series | [OpenMP,5.0] Implement relaxation of implicit map vs. existing device mappings (for mainline trunk) | expand |
Hi Chung-Lin! On 2021-05-14T21:20:25+0800, Chung-Lin Tang <cltang@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, <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 <cltang@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
On Fri, May 14, 2021 at 09:20:25PM +0800, Chung-Lin Tang wrote: > 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) > { As Thomas mentioned, there is now also OMP_CLAUSE_MAP_IMPLICIT that means something different: /* 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. */ Having OMP_CLAUSE_MAP_IMPLICIT and OMP_CLAUSE_MAP_IMPLICIT_P would be too confusing. So either we need to use just one flag for both purposes or have two different flags and find a better name for one of them. The former would be possible if no OMP_CLAUSE_MAP clauses added by the FEs are implicit - then you could clear OMP_CLAUSE_MAP_IMPLICIT in gimplify_scan_omp_clauses. I wonder if it is the case though, e.g. doesn't your "Improve OpenMP target support for C++ [PR92120 v4]" patch add a lot of such implicit map clauses (e.g. the this[:1] and various others)? Also, gimplify_adjust_omp_clauses_1 sometimes doesn't add just one map clause, but several, shouldn't those be marked implicit too? And similarly it calls lang_hooks.decls.omp_finish_clause which can add even further map clauses implicitly, shouldn't those be implicit too (in that case copy the flag from the clause it is called on to the extra clauses it adds)? Also as Thomas mentioned, it should be restricted to non-OpenACC, it can check gimplify_omp_ctxp->region_type if it is OpenMP or OpenACC. > @@ -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 Two spaces after dot in comments. > + 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); Why are the implicit map clauses added first and not last? There is also the OpenMP 5.1 [352:17-22] case which basically says that the implicit mappings should be ignored if there are explicit ones on the same construct (though, do we really create implicit clauses in that case?). > --- 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) I think here we need to decide with which GOMP_MAP* kinds the implicit bit will need to be combined with, with looking forward into what features we still need to implement for OpenMP 5.0/5.1 (not aware of anything in 5.2 that would need special care but perhaps I've missed it). E.g. for declare mapper those single OMP_CLAUSE_MAPs with the implicit bit might need to be split into various smaller ones, map this FIELD_DECL, map that other FIELD_DECL, what it points to, etc. Even without declare mapper, the spec now says that mapping a structure is as if each member is mapped separately and 5.2 is going to say there is a predefined declare mapper that does that (which is of course something we don't want under the hood, we don't want thousands of maps, but virtually split it on a field by field basis, do all the special stuff - e.g. how pointers are to be mapped as zero array sections with pointer attachment, how C++ references are to be handled etc. and then virtually coalesce all the adjacent fields that have the same treatment back). I think that actually means we should defer most of the map struct etc. handling we do in gimplify_scan_omp_clauses, instead note the explicit map clauses and what they refer to, add implicit ones (and set implicit bit on those), then go over all clauses (explicit and implicit), do the declare mapper processing, do the sorting of what that produces based on base expressions/pointers etc. At this point, I'm not sure if GOMP_MAP_IMPLICIT can or can't appear together e.g. with GOMP_MAP_STRUCT. > @@ -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; As the particular bit isn't used for anything right now, perhaps just do val &= ~GOMP_MAP_IMPLICIT unconditionally? But, on the other side, for !short_mapkind you do not want to mask that bit out, only the low 3 bits are the mapping type and the upper bits are log2 of alignment, so the above for !short_mapkind would in weird way change some alignments. > + 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); Similarly can return 0 for !short_mapkind, the compatibility GOMP_target etc. APIs will never have those in. Jakub
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; +}
Hi Jakub, This is a version of patch https://gcc.gnu.org/pipermail/gcc-patches/2021-May/569665.html for mainline trunk. 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 <cltang@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.