* [PATCH, OG10, OpenMP 5.0, committed] Implement relaxation of implicit map vs. existing device mappings
@ 2021-05-05 15:17 Chung-Lin Tang
2021-05-07 12:35 ` Thomas Schwinge
0 siblings, 1 reply; 9+ messages in thread
From: Chung-Lin Tang @ 2021-05-05 15:17 UTC (permalink / raw)
To: gcc-patches, Catherine Moore, Tobias Burnus
[-- Attachment #1: Type: text/plain, Size: 5989 bytes --]
This patch implements relaxing the requirements when a map with the implicit attribute encounters
an overlapping existing map. As the OpenMP 5.0 spec describes on page 320, lines 18-27 (and 5.1 spec,
page 352, lines 13-22):
"If a single contiguous part of the original storage of a list item with an implicit data-mapping
attribute has corresponding storage in the device data environment prior to a task encountering the
construct that is associated with the map clause, only that part of the original storage will have
corresponding storage in the device data environment as a result of the map clause."
Also tracked in the OpenMP spec context as issue #1463:
https://github.com/OpenMP/spec/issues/1463
The implementation inside the compiler is to of course, tag the implicitly created maps with some
indication of "implicit". I've done this with a OMP_CLAUSE_MAP_IMPLICIT_P macro, using
'base.deprecated_flag' underneath.
There is an encoding of this as GOMP_MAP_IMPLICIT == GOMP_MAP_FLAG_SPECIAL_3|GOMP_MAP_FLAG_SPECIAL_4
in include/gomp-constants.h for the runtime, but I've intentionally avoided exploding the entire
gimplify/omp-low with a new set of GOMP_MAP_IMPLICIT_TO/FROM/etc. symbols, instead adding in the new
flag bits only at the final runtime call generation during omp-lowering.
The rest is libgomp mapping taking care of the implicit case: allowing map success if an existing
map is a proper subset of the new map, if the new map is implicit. Straightforward enough I think.
There are also some additions to print the implicit attribute during tree pretty-printing, for that
reason some scan tests were updated.
Also, another adjustment in this patch is how implicitly created clauses are added to the current
clause list in gimplify_adjust_omp_clauses(). Instead of simply appending the new clauses to the end,
this patch adds them at the position "after initial non-map clauses, but right before any existing
map clauses".
The reason for this is: when combined with other map clauses, for example:
#pragma omp target map(rec.ptr[:N])
for (int i = 0; i < N; i++)
rec.ptr[i] += 1;
There will be an implicit map created for map(rec), because of the access inside the target region.
The expectation is that 'rec' is implicitly mapped, and then the pointed array-section part by 'rec.ptr'
will be mapped, and then attachment to the 'rec.ptr' field of the mapped 'rec' (in that order).
If the implicit 'map(rec)' is appended to the end, instead of placed before other maps, the attachment
operation will not find anything to attach to, and the entire region will fail.
Note: this touches a bit on another issue which I will be sending a patch for later:
per the discussion on omp-lang, an array section list item should *not* be mapping its base-pointer
(although an attachment attempt should exist), while in current GCC behavior, for struct member pointers
like 'rec.ptr' above, we do map it (which should be deemed incorrect).
This means that as of right now, this modification of map order doesn't really exhibit the above mentioned
behavior yet. I have included it as part of this patch because the "[implicit]" tree printing requires
modifying many gimple scan tests already, so including the test modifications together seems more
manageable patch-wise.
Tested with no regressions, and pushed to devel/omp/gcc-10. Will be submitting a mainline trunk version later.
Chung-Lin
2021-05-05 Chung-Lin Tang <cltang@codesourcery.com>
include/ChangeLog:
* gomp-constants.h (GOMP_MAP_IMPLICIT): New special map kind bits value.
(GOMP_MAP_FLAG_SPECIAL_BITS): Define helper mask for whole set of
special map kind bits.
(GOMP_MAP_NONCONTIG_ARRAY_P): Adjust test for non-contiguous array map
kind bits to be more specific.
(GOMP_MAP_IMPLICIT_P): New predicate macro for implicit map kinds.
gcc/ChangeLog:
* tree.h (OMP_CLAUSE_MAP_IMPLICIT_P): New access macro for 'implicit'
bit, using 'base.deprecated_flag' field of tree_node.
* tree-pretty-print.c (dump_omp_clause): Add support for printing
implicit attribute in tree dumping.
* gimplify.c (gimplify_adjust_omp_clauses_1):
Set OMP_CLAUSE_MAP_IMPLICIT_P to 1 if map clause is implicitly created.
(gimplify_adjust_omp_clauses): Adjust place of adding implicitly created
clauses, from simple append, to starting of list, after non-map clauses.
* omp-low.c (lower_omp_target): Add GOMP_MAP_IMPLICIT bits into kind
values passed to libgomp for implicit maps.
gcc/testsuite/ChangeLog:
* c-c++-common/gomp/target-implicit-map-1.c: New test.
* c-c++-common/goacc/combined-reduction.c: Adjust scan test pattern.
* c-c++-common/goacc/firstprivate-mappings-1.c: Likewise.
* c-c++-common/goacc/mdc-1.c: Likewise.
* c-c++-common/goacc/reduction-1.c: Likewise.
* c-c++-common/goacc/reduction-2.c: Likewise.
* c-c++-common/goacc/reduction-3.c: Likewise.
* c-c++-common/goacc/reduction-4.c: Likewise.
* c-c++-common/goacc/reduction-8.c: Likewise.
* g++.dg/goacc/firstprivate-mappings-1.C: Likewise.
* g++.dg/gomp/target-lambda-1.C: Likewise.
* g++.dg/gomp/target-this-3.C: Likewise.
* g++.dg/gomp/target-this-4.C: Likewise.
* gfortran.dg/goacc/common-block-3.f90: Likewise.
* gfortran.dg/goacc/loop-tree-1.f90: Likewise.
* gfortran.dg/goacc/private-explicit-kernels-1.f95: Likewise.
* gfortran.dg/goacc/private-predetermined-kernels-1.f95: Likewise.
libgomp/ChangeLog:
* target.c (gomp_map_vars_existing): Add 'bool implicit' parameter, add
implicit map handling to allow a "superset" existing map as valid case.
(get_kind): Adjust to filter out GOMP_MAP_IMPLICIT bits in return value.
(get_implicit): New function to extract implicit status.
(gomp_map_fields_existing): Adjust arguments in calls to
gomp_map_vars_existing, and add uses of get_implicit.
(gomp_map_vars_internal): Likewise.
* testsuite/libgomp.c-c++-common/target-implicit-map-1.c: New test.
[-- Attachment #2: 0001-OpenMP-5.0-Implement-relaxation-of-implicit-map-vs.-.patch --]
[-- Type: text/plain, Size: 41776 bytes --]
From a70b5b1aa8b3d32f6728dbfcfc00b0cff8c5219d Mon Sep 17 00:00:00 2001
From: Chung-Lin Tang <cltang@codesourcery.com>
Date: Wed, 5 May 2021 08:11:19 -0700
Subject: [PATCH] OpenMP 5.0: Implement relaxation of implicit map vs. existing
device mappings
This patch implements relaxing the requirements when a map with the implicit
attribute encounters an overlapping existing map. As the OpenMP 5.0 spec
describes on page 320, lines 18-27 (and 5.1 spec, page 352, lines 13-22):
"If a single contiguous part of the original storage of a list item with an
implicit data-mapping attribute has corresponding storage in the device data
environment prior to a task encountering the construct that is associated with
the map clause, only that part of the original storage will have corresponding
storage in the device data environment as a result of the map clause."
Also tracked in the OpenMP spec context as issue #1463:
https://github.com/OpenMP/spec/issues/1463
2021-05-05 Chung-Lin Tang <cltang@codesourcery.com>
include/ChangeLog:
* gomp-constants.h (GOMP_MAP_IMPLICIT): New special map kind bits value.
(GOMP_MAP_FLAG_SPECIAL_BITS): Define helper mask for whole set of
special map kind bits.
(GOMP_MAP_NONCONTIG_ARRAY_P): Adjust test for non-contiguous array map
kind bits to be more specific.
(GOMP_MAP_IMPLICIT_P): New predicate macro for implicit map kinds.
gcc/ChangeLog:
* tree.h (OMP_CLAUSE_MAP_IMPLICIT_P): New access macro for 'implicit'
bit, using 'base.deprecated_flag' field of tree_node.
* tree-pretty-print.c (dump_omp_clause): Add support for printing
implicit attribute in tree dumping.
* gimplify.c (gimplify_adjust_omp_clauses_1):
Set OMP_CLAUSE_MAP_IMPLICIT_P to 1 if map clause is implicitly created.
(gimplify_adjust_omp_clauses): Adjust place of adding implicitly created
clauses, from simple append, to starting of list, after non-map clauses.
* omp-low.c (lower_omp_target): Add GOMP_MAP_IMPLICIT bits into kind
values passed to libgomp for implicit maps.
gcc/testsuite/ChangeLog:
* c-c++-common/gomp/target-implicit-map-1.c: New test.
* c-c++-common/goacc/combined-reduction.c: Adjust scan test pattern.
* c-c++-common/goacc/firstprivate-mappings-1.c: Likewise.
* c-c++-common/goacc/mdc-1.c: Likewise.
* c-c++-common/goacc/reduction-1.c: Likewise.
* c-c++-common/goacc/reduction-2.c: Likewise.
* c-c++-common/goacc/reduction-3.c: Likewise.
* c-c++-common/goacc/reduction-4.c: Likewise.
* c-c++-common/goacc/reduction-8.c: Likewise.
* g++.dg/goacc/firstprivate-mappings-1.C: Likewise.
* g++.dg/gomp/target-lambda-1.C: Likewise.
* g++.dg/gomp/target-this-3.C: Likewise.
* g++.dg/gomp/target-this-4.C: Likewise.
* gfortran.dg/goacc/common-block-3.f90: Likewise.
* gfortran.dg/goacc/loop-tree-1.f90: Likewise.
* gfortran.dg/goacc/private-explicit-kernels-1.f95: Likewise.
* gfortran.dg/goacc/private-predetermined-kernels-1.f95: Likewise.
libgomp/ChangeLog:
* target.c (gomp_map_vars_existing): Add 'bool implicit' parameter, add
implicit map handling to allow a "superset" existing map as valid case.
(get_kind): Adjust to filter out GOMP_MAP_IMPLICIT bits in return value.
(get_implicit): New function to extract implicit status.
(gomp_map_fields_existing): Adjust arguments in calls to
gomp_map_vars_existing, and add uses of get_implicit.
(gomp_map_vars_internal): Likewise.
* testsuite/libgomp.c-c++-common/target-implicit-map-1.c: New test.
---
gcc/gimplify.c | 11 ++-
gcc/omp-low.c | 13 ++++
.../c-c++-common/goacc/combined-reduction.c | 4 +-
.../c-c++-common/goacc/firstprivate-mappings-1.c | 6 +-
gcc/testsuite/c-c++-common/goacc/mdc-1.c | 2 +-
gcc/testsuite/c-c++-common/goacc/reduction-1.c | 4 +-
gcc/testsuite/c-c++-common/goacc/reduction-2.c | 4 +-
gcc/testsuite/c-c++-common/goacc/reduction-3.c | 4 +-
gcc/testsuite/c-c++-common/goacc/reduction-4.c | 4 +-
gcc/testsuite/c-c++-common/goacc/reduction-8.c | 10 +--
.../c-c++-common/gomp/target-implicit-map-1.c | 39 +++++++++++
.../g++.dg/goacc/firstprivate-mappings-1.C | 2 +-
gcc/testsuite/g++.dg/gomp/target-lambda-1.C | 6 +-
gcc/testsuite/g++.dg/gomp/target-this-3.C | 4 +-
gcc/testsuite/g++.dg/gomp/target-this-4.C | 4 +-
gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 | 8 +--
gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90 | 2 +-
.../goacc/private-explicit-kernels-1.f95 | 4 +-
.../goacc/private-predetermined-kernels-1.f95 | 4 +-
gcc/tree-pretty-print.c | 3 +
gcc/tree.h | 4 ++
include/gomp-constants.h | 17 ++++-
libgomp/target.c | 78 ++++++++++++++++------
.../libgomp.c-c++-common/target-implicit-map-1.c | 31 +++++++++
24 files changed, 211 insertions(+), 57 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/gomp/target-implicit-map-1.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-1.c
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 91aa15d..ba071e8 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -10579,6 +10579,7 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
gcc_unreachable ();
}
OMP_CLAUSE_SET_MAP_KIND (clause, kind);
+ OMP_CLAUSE_MAP_IMPLICIT_P (clause) = 1;
if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
{
@@ -11158,9 +11159,15 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
list_p = &OMP_CLAUSE_CHAIN (c);
}
- /* Add in any implicit data sharing. */
+ /* Add in any implicit data sharing. Implicit clauses are added at the start
+ of the clause list, but after any non-map clauses. */
struct gimplify_adjust_omp_clauses_data data;
- data.list_p = list_p;
+ tree *implicit_add_list_p = orig_list_p;
+ while (*implicit_add_list_p
+ && OMP_CLAUSE_CODE (*implicit_add_list_p) != OMP_CLAUSE_MAP)
+ implicit_add_list_p = &OMP_CLAUSE_CHAIN (*implicit_add_list_p);
+
+ data.list_p = implicit_add_list_p;
data.pre_p = pre_p;
splay_tree_foreach (ctx->variables, gimplify_adjust_omp_clauses_1, &data);
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 66519ad..64b7c19 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -12920,6 +12920,19 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
else if (integer_nonzerop (s))
tkind_zero = tkind;
}
+ if (tkind_zero == tkind
+ && OMP_CLAUSE_MAP_IMPLICIT_P (c)
+ && (((tkind & GOMP_MAP_FLAG_SPECIAL_BITS)
+ & ~GOMP_MAP_IMPLICIT)
+ == 0))
+ {
+ /* If this is an implicit map, and the GOMP_MAP_IMPLICIT
+ bits are not interfered by other special bit encodings,
+ then turn the GOMP_IMPLICIT_BIT flag on for the runtime
+ to see. */
+ tkind |= GOMP_MAP_IMPLICIT;
+ tkind_zero = tkind;
+ }
break;
case OMP_CLAUSE_FIRSTPRIVATE:
gcc_checking_assert (is_gimple_omp_oacc (ctx->stmt));
diff --git a/gcc/testsuite/c-c++-common/goacc/combined-reduction.c b/gcc/testsuite/c-c++-common/goacc/combined-reduction.c
index ecf23f5..fa67e08 100644
--- a/gcc/testsuite/c-c++-common/goacc/combined-reduction.c
+++ b/gcc/testsuite/c-c++-common/goacc/combined-reduction.c
@@ -23,7 +23,7 @@ main ()
return 0;
}
-/* { dg-final { scan-tree-dump-times "omp target oacc_parallel reduction.+:v1. map.tofrom:v1" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "omp target oacc_parallel reduction.+:v1. firstprivate.n. map.tofrom:v1" 1 "gimple" } } */
/* { dg-final { scan-tree-dump-times "acc loop reduction.+:v1. private.i." 1 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "omp target oacc_kernels map.force_tofrom:n .len: 4.. map.force_tofrom:v1 .len: 4.." 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "omp target oacc_kernels map.force_tofrom:n .len: 4..implicit.. map.force_tofrom:v1 .len: 4..implicit.." 1 "gimple" } } */
/* { dg-final { scan-tree-dump-times "acc loop reduction.+:v1. private.i." 1 "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c b/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c
index 7987bea..f43e4b4 100644
--- a/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c
@@ -419,12 +419,12 @@ vla (int array_li)
copyout (array_so)
/* The gimplifier has created an implicit 'firstprivate' clause for the array
length.
- { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel map\(from:array_so \[len: 4\]\) firstprivate\(array_li.[0-9]+\)} omplower { target { ! c++ } } } }
- { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel map\(from:array_so \[len: 4\]\) firstprivate\(} omplower { target { c++ } } } }
+ { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel firstprivate\(array_li.[0-9]+\) map\(from:array_so \[len: 4\]\)} omplower { target { ! c++ } } } }
+ { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel firstprivate\([^)]+\) map\(from:array_so \[len: 4\]\)} omplower { target { c++ } } } }
(C++ computes an intermediate value, so can't scan for 'firstprivate(array_li)'.) */
/* For C, non-LP64, the gimplifier has also created a mapping for the array
itself; PR90859.
- { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel map\(from:array_so \[len: 4\]\) firstprivate\(array_li.[0-9]+\) map\(tofrom:\(\*array.[0-9]+\) \[len: D\.[0-9]+\]\) map\(firstprivate:array \[pointer assign, bias: 0\]\) \[} omplower { target { c && { ! lp64 } } } } } */
+ { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel firstprivate\(array_li.[0-9]+\) map\(tofrom:\(\*array.[0-9]+\) \[len: D\.[0-9]+\]\[implicit\]\) map\(firstprivate:array \[pointer assign, bias: 0\]\) map\(from:array_so \[len: 4\]\) \[} omplower { target { c && { ! lp64 } } } } } */
{
array_so = sizeof array;
}
diff --git a/gcc/testsuite/c-c++-common/goacc/mdc-1.c b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
index 337c1f7..9f43de4 100644
--- a/gcc/testsuite/c-c++-common/goacc/mdc-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
@@ -45,7 +45,7 @@ t1 ()
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:s .len: 32.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.tofrom:.z .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len: 8.. map.tofrom:._1 .len: 40.. map.attach:s.a .bias: 0.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.attach:s.e .bias: 0.. map.tofrom:s .len: 32" 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.tofrom:s .len: 32..implicit.. map.attach:s.e .bias: 0.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.attach:a .bias: 0.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:a .bias: 0.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:a .len: 8.." 1 "omplower" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-1.c b/gcc/testsuite/c-c++-common/goacc/reduction-1.c
index 35bfc86..d9e3c38 100644
--- a/gcc/testsuite/c-c++-common/goacc/reduction-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/reduction-1.c
@@ -68,5 +68,5 @@ main(void)
}
/* Check that default copy maps are generated for loop reductions. */
-/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\]\\)" 7 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 7 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 2 "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-2.c b/gcc/testsuite/c-c++-common/goacc/reduction-2.c
index 9dba035..18dc03c 100644
--- a/gcc/testsuite/c-c++-common/goacc/reduction-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/reduction-2.c
@@ -50,5 +50,5 @@ main(void)
}
/* Check that default copy maps are generated for loop reductions. */
-/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\]\\)" 4 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 4 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 2 "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-3.c b/gcc/testsuite/c-c++-common/goacc/reduction-3.c
index 669cd43..2311d4b 100644
--- a/gcc/testsuite/c-c++-common/goacc/reduction-3.c
+++ b/gcc/testsuite/c-c++-common/goacc/reduction-3.c
@@ -50,5 +50,5 @@ main(void)
}
/* Check that default copy maps are generated for loop reductions. */
-/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\]\\)" 4 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 4 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 2 "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-4.c b/gcc/testsuite/c-c++-common/goacc/reduction-4.c
index 5c3dfb1..57823f8 100644
--- a/gcc/testsuite/c-c++-common/goacc/reduction-4.c
+++ b/gcc/testsuite/c-c++-common/goacc/reduction-4.c
@@ -38,5 +38,5 @@ main(void)
}
/* Check that default copy maps are generated for loop reductions. */
-/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 2 "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-8.c b/gcc/testsuite/c-c++-common/goacc/reduction-8.c
index 8a0283f..8494e59 100644
--- a/gcc/testsuite/c-c++-common/goacc/reduction-8.c
+++ b/gcc/testsuite/c-c++-common/goacc/reduction-8.c
@@ -87,8 +87,10 @@ main(void)
/* Check that default copy maps are generated for loop reductions. */
/* { dg-final { scan-tree-dump-times "reduction..:result. map.tofrom:result .len: 4.." 1 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "oacc_parallel map.tofrom:result .len: 4.." 2 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000.. firstprivate.result." 3 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "map.tofrom:result .len: 4.. map.tofrom:array .len: 4000.." 1 "gimple" } } */
-/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000.. map.force_tofrom:result .len: 4.." 1 "gimple" } } */
+
+/* { dg-final { scan-tree-dump-times "oacc_parallel map.tofrom:result .len: 4..implicit.." 1 "gimple" } } */
+
+/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000..implicit.. firstprivate.result." 3 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000..implicit.. map.tofrom:result .len: 4.." 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000..implicit.. map.force_tofrom:result .len: 4..implicit.." 1 "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/target-implicit-map-1.c b/gcc/testsuite/c-c++-common/gomp/target-implicit-map-1.c
new file mode 100644
index 0000000..52944fd
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-implicit-map-1.c
@@ -0,0 +1,39 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+#ifdef __cplusplus
+extern "C"
+#else
+extern
+#endif
+void abort (void);
+
+int
+main (void)
+{
+ #define N 5
+ int array[N][N];
+
+ for (int i = 0; i < N; i++)
+ {
+ #pragma omp target enter data map(alloc: array[i:1][0:N])
+
+ #pragma omp target
+ for (int j = 0; j < N; j++)
+ array[i][j] = i * 10 + j;
+
+ #pragma omp target exit data map(from: array[i:1][0:N])
+ }
+
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < N; j++)
+ if (array[i][j] != i + j)
+ abort ();
+
+ return 0;
+}
+
+/* { dg-final { scan-tree-dump {#pragma omp target enter data map\(alloc:array\[[^]]+\]\[0\] \[len: [0-9]+\]\)} "gimple" } } */
+
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(i\) map\(tofrom:array \[len: [0-9]+\]\[implicit\]\)} "gimple" } } */
+
+/* { dg-final { scan-tree-dump {#pragma omp target exit data map\(from:array\[[^]]+\]\[0\] \[len: [0-9]+\]\)} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C b/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C
index 1b1badb..99a3bd4 100644
--- a/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C
+++ b/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C
@@ -416,7 +416,7 @@ vla (int &array_li)
copyout (array_so)
/* The gimplifier has created an implicit 'firstprivate' clause for the array
length.
- { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel map\(from:array_so \[len: 4\]\) firstprivate\(} omplower } }
+ { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel firstprivate\([^)]+\) map\(from:array_so \[len: 4\]\)} omplower } }
(C++ computes an intermediate value, so can't scan for 'firstprivate(array_li)'.) */
{
array_so = sizeof array;
diff --git a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
index 7dceef8..e5a24d7 100644
--- a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
+++ b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
@@ -87,8 +87,8 @@ int main (void)
return 0;
}
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(end\) firstprivate\(begin\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data2 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(end\) firstprivate\(begin\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data2 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-3.C b/gcc/testsuite/g++.dg/gomp/target-this-3.C
index 08568f9..2755b4b 100644
--- a/gcc/testsuite/g++.dg/gomp/target-this-3.C
+++ b/gcc/testsuite/g++.dg/gomp/target-this-3.C
@@ -100,6 +100,6 @@ int main (void)
return 0;
}
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:this->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9+] \[len: 0\]\) firstprivate\(n\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:this->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9+] \[len: 0\]\)} "gimple" } } */
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:this->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(n\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:this->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-4.C b/gcc/testsuite/g++.dg/gomp/target-this-4.C
index 3b2d581..3703762 100644
--- a/gcc/testsuite/g++.dg/gomp/target-this-4.C
+++ b/gcc/testsuite/g++.dg/gomp/target-this-4.C
@@ -102,6 +102,6 @@ int main (void)
return 0;
}
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: 1\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: 1\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:_[0-9]+->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:_[0-9]+->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */
diff --git a/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 b/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90
index e43d376..e9f169f 100644
--- a/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90
@@ -33,10 +33,10 @@ end program main
! { dg-final { scan-tree-dump-times "omp target oacc_data_kernels .*map\\(tofrom:x \\\[len: 400\\\]\\)" 1 "omplower" } }
! { dg-final { scan-tree-dump-times "omp target oacc_data_kernels .*map\\(tofrom:y \\\[len: 400\\\]\\\)" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:x \\\[len: 400\\\]\\)" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:y \\\[len: 400\\\]\\\)" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_tofrom:i \\\[len: 4\\\]\\)" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_tofrom:c \\\[len: 4\\\]\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:x \\\[len: 400\\\]\\\[implicit\\\]\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:y \\\[len: 400\\\]\\\[implicit\\\]\\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_tofrom:i \\\[len: 4\\\]\\\[implicit\\\]\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_tofrom:c \\\[len: 4\\\]\\\[implicit\\\]\\)" 1 "omplower" } }
! Expecting no mapping of un-referenced common-blocks variables
diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90 b/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90
index 150f930..4cdfc55 100644
--- a/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90
@@ -44,4 +44,4 @@ end program test
! { dg-final { scan-tree-dump-times "private\\(m\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "reduction\\(\\+:sum\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "map\\(tofrom:sum \\\[len: \[0-9\]+\\\]\\)" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map\\(tofrom:sum \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 1 "gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95 b/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95
index 0c47045..fef5126 100644
--- a/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95
@@ -83,7 +83,7 @@ program test
!$acc kernels ! Explicit "private(i2_2_s)" clause cannot be specified here.
! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "original" { xfail *-*-* } } } ! PR90067
! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
- ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i2_2_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i2_2_s \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
do i2_2_s = 1, 100
!$acc loop private(j2_2_s) independent
! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_2_s\\) independent" 1 "original" } }
@@ -234,7 +234,7 @@ program test
!$acc kernels ! Explicit "private(i3_5_s)" clause cannot be specified here.
! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "original" { xfail *-*-* } } } ! PR90067
! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
- ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i3_5_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i3_5_s \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
do i3_5_s = 1, 100
!$acc loop private(j3_5_s) independent
! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_5_s\\) independent" 1 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95 b/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95
index 3357a20..38459cf 100644
--- a/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95
@@ -83,7 +83,7 @@ program test
!$acc kernels
! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "original" { xfail *-*-* } } } ! PR90067
! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
- ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i2_2_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i2_2_s \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
do i2_2_s = 1, 100
!$acc loop independent
! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_2_s\\) independent" 1 "original" } }
@@ -234,7 +234,7 @@ program test
!$acc kernels
! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "original" { xfail *-*-* } } } ! PR90067
! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
- ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i3_5_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
+ ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i3_5_s \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
do i3_5_s = 1, 100
!$acc loop independent
! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_5_s\\) independent" 1 "original" } }
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 261cc9d..4cd4606 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -946,6 +946,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
spc, flags, false);
pp_right_bracket (pp);
}
+ if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_IMPLICIT_P (clause))
+ pp_string (pp, "[implicit]");
pp_right_paren (pp);
break;
diff --git a/gcc/tree.h b/gcc/tree.h
index 8d9829c..647c5ba 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1637,6 +1637,10 @@ class auto_suppress_location_wrappers
variable. */
#define OMP_CLAUSE_MAP_IN_REDUCTION(NODE) \
TREE_PRIVATE (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
+/* Nonzero if this map clause was created through implicit data-mapping
+ rules. */
+#define OMP_CLAUSE_MAP_IMPLICIT_P(NODE) \
+ (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.deprecated_flag)
/* True on an OMP_CLAUSE_USE_DEVICE_PTR with an OpenACC 'if_present'
clause. */
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index b8efb30..33cfcb9 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -46,6 +46,16 @@
| GOMP_MAP_FLAG_SPECIAL_0)
#define GOMP_MAP_DEEP_COPY (GOMP_MAP_FLAG_SPECIAL_4 \
| GOMP_MAP_FLAG_SPECIAL_2)
+/* This value indicates the map was created implicitly according to
+ OpenMP rules. */
+#define GOMP_MAP_IMPLICIT (GOMP_MAP_FLAG_SPECIAL_3 \
+ | GOMP_MAP_FLAG_SPECIAL_4)
+/* Mask for entire set of special map kind bits. */
+#define GOMP_MAP_FLAG_SPECIAL_BITS (GOMP_MAP_FLAG_SPECIAL_0 \
+ | GOMP_MAP_FLAG_SPECIAL_1 \
+ | GOMP_MAP_FLAG_SPECIAL_2 \
+ | GOMP_MAP_FLAG_SPECIAL_3 \
+ | GOMP_MAP_FLAG_SPECIAL_4)
/* Flag to force a specific behavior (or else, trigger a run-time error). */
#define GOMP_MAP_FLAG_FORCE (1 << 7)
@@ -225,7 +235,12 @@ enum gomp_map_kind
(GOMP_MAP_ALWAYS_TO_P (X) || ((X) == GOMP_MAP_ALWAYS_FROM))
#define GOMP_MAP_NONCONTIG_ARRAY_P(X) \
- ((X) & GOMP_MAP_NONCONTIG_ARRAY)
+ (((X) & GOMP_MAP_FLAG_SPECIAL_BITS) == GOMP_MAP_NONCONTIG_ARRAY \
+ || (X) == GOMP_MAP_NONCONTIG_ARRAY_FORCE_PRESENT)
+
+#define GOMP_MAP_IMPLICIT_P(X) \
+ (((X) & GOMP_MAP_FLAG_SPECIAL_BITS) == GOMP_MAP_IMPLICIT)
+
/* Asynchronous behavior. Keep in sync with
libgomp/{openacc.h,openacc.f90,openacc_lib.h}:acc_async_t. */
diff --git a/libgomp/target.c b/libgomp/target.c
index 9c75826..ecda2ef 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -510,7 +510,7 @@ static inline void
gomp_map_vars_existing (struct gomp_device_descr *devicep,
struct goacc_asyncqueue *aq, splay_tree_key oldn,
splay_tree_key newn, struct target_var_desc *tgt_var,
- unsigned char kind, bool always_to_flag,
+ unsigned char kind, bool always_to_flag, bool implicit,
struct gomp_coalesce_buf *cbuf,
htab_t *refcount_set)
{
@@ -522,11 +522,22 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep,
tgt_var->always_copy_from = GOMP_MAP_ALWAYS_FROM_P (kind);
tgt_var->is_attach = false;
tgt_var->offset = newn->host_start - oldn->host_start;
- tgt_var->length = newn->host_end - newn->host_start;
+
+ /* For implicit maps, old contained in new is valid. */
+ bool implicit_subset = (implicit
+ && newn->host_start <= oldn->host_start
+ && oldn->host_end <= newn->host_end);
+ if (implicit_subset)
+ tgt_var->length = oldn->host_end - oldn->host_start;
+ else
+ tgt_var->length = newn->host_end - newn->host_start;
if ((kind & GOMP_MAP_FLAG_FORCE)
- || oldn->host_start > newn->host_start
- || oldn->host_end < newn->host_end)
+ /* For implicit maps, old contained in new is valid. */
+ || !(implicit_subset
+ /* Otherwise, new contained inside old is considered valid. */
+ || (oldn->host_start <= newn->host_start
+ && newn->host_end <= oldn->host_end)))
{
gomp_mutex_unlock (&devicep->lock);
gomp_fatal ("Trying to map into device [%p..%p) object when "
@@ -536,11 +547,17 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep,
}
if (GOMP_MAP_ALWAYS_TO_P (kind) || always_to_flag)
- gomp_copy_host2dev (devicep, aq,
- (void *) (oldn->tgt->tgt_start + oldn->tgt_offset
- + newn->host_start - oldn->host_start),
- (void *) newn->host_start,
- newn->host_end - newn->host_start, false, cbuf);
+ {
+ /* Implicit + always should not happen. If this does occur, below
+ address/length adjustment is a TODO. */
+ assert (!implicit_subset);
+
+ gomp_copy_host2dev (devicep, aq,
+ (void *) (oldn->tgt->tgt_start + oldn->tgt_offset
+ + newn->host_start - oldn->host_start),
+ (void *) newn->host_start,
+ newn->host_end - newn->host_start, false, cbuf);
+ }
gomp_increment_refcount (oldn, refcount_set);
}
@@ -548,8 +565,24 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep,
static int
get_kind (bool short_mapkind, void *kinds, int idx)
{
- return short_mapkind ? ((unsigned short *) kinds)[idx]
- : ((unsigned char *) kinds)[idx];
+ int val = (short_mapkind
+ ? ((unsigned short *) kinds)[idx]
+ : ((unsigned char *) kinds)[idx]);
+
+ if (GOMP_MAP_IMPLICIT_P (val))
+ val &= ~GOMP_MAP_IMPLICIT;
+ return val;
+}
+
+
+static bool
+get_implicit (bool short_mapkind, void *kinds, int idx)
+{
+ int val = (short_mapkind
+ ? ((unsigned short *) kinds)[idx]
+ : ((unsigned char *) kinds)[idx]);
+
+ return GOMP_MAP_IMPLICIT_P (val);
}
static void
@@ -612,6 +645,7 @@ gomp_map_fields_existing (struct target_mem_desc *tgt,
struct splay_tree_s *mem_map = &devicep->mem_map;
struct splay_tree_key_s cur_node;
int kind;
+ bool implicit;
const bool short_mapkind = true;
const int typemask = short_mapkind ? 0xff : 0x7;
@@ -619,12 +653,14 @@ gomp_map_fields_existing (struct target_mem_desc *tgt,
cur_node.host_end = cur_node.host_start + sizes[i];
splay_tree_key n2 = splay_tree_lookup (mem_map, &cur_node);
kind = get_kind (short_mapkind, kinds, i);
+ implicit = get_implicit (short_mapkind, kinds, i);
if (n2
&& n2->tgt == n->tgt
&& n2->host_start - n->host_start == n2->tgt_offset - n->tgt_offset)
{
gomp_map_vars_existing (devicep, aq, n2, &cur_node, &tgt->list[i],
- kind & typemask, false, cbuf, refcount_set);
+ kind & typemask, false, implicit, cbuf,
+ refcount_set);
return;
}
if (sizes[i] == 0)
@@ -640,7 +676,8 @@ gomp_map_fields_existing (struct target_mem_desc *tgt,
== n2->tgt_offset - n->tgt_offset)
{
gomp_map_vars_existing (devicep, aq, n2, &cur_node, &tgt->list[i],
- kind & typemask, false, cbuf, refcount_set);
+ kind & typemask, false, implicit, cbuf,
+ refcount_set);
return;
}
}
@@ -652,7 +689,8 @@ gomp_map_fields_existing (struct target_mem_desc *tgt,
&& n2->host_start - n->host_start == n2->tgt_offset - n->tgt_offset)
{
gomp_map_vars_existing (devicep, aq, n2, &cur_node, &tgt->list[i],
- kind & typemask, false, cbuf, refcount_set);
+ kind & typemask, false, implicit, cbuf,
+ refcount_set);
return;
}
}
@@ -898,6 +936,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
for (i = 0; i < mapnum; i++)
{
int kind = get_kind (short_mapkind, kinds, i);
+ bool implicit = get_implicit (short_mapkind, kinds, i);
if (hostaddrs[i] == NULL
|| (kind & typemask) == GOMP_MAP_FIRSTPRIVATE_INT)
{
@@ -1104,8 +1143,8 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
}
}
gomp_map_vars_existing (devicep, aq, n, &cur_node, &tgt->list[i],
- kind & typemask, always_to_cnt > 0, NULL,
- refcount_set);
+ kind & typemask, always_to_cnt > 0, implicit,
+ NULL, refcount_set);
i += always_to_cnt;
}
else
@@ -1182,7 +1221,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
{
assert (n->refcount != REFCOUNT_LINK);
gomp_map_vars_existing (devicep, aq, n, &cur_node, row_desc,
- kind & typemask, false,
+ kind & typemask, false, false,
/* TODO: cbuf? */ NULL, refcount_set);
}
else
@@ -1312,6 +1351,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
else if (tgt->list[i].key == NULL)
{
int kind = get_kind (short_mapkind, kinds, i);
+ bool implicit = get_implicit (short_mapkind, kinds, i);
if (hostaddrs[i] == NULL)
continue;
switch (kind & typemask)
@@ -1483,7 +1523,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
splay_tree_key n = splay_tree_lookup (mem_map, k);
if (n && n->refcount != REFCOUNT_LINK)
gomp_map_vars_existing (devicep, aq, n, k, &tgt->list[i],
- kind & typemask, false, cbufp,
+ kind & typemask, false, implicit, cbufp,
refcount_set);
else
{
@@ -1702,7 +1742,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
{
assert (k->refcount != REFCOUNT_LINK);
gomp_map_vars_existing (devicep, aq, k, &cur_node, row_desc,
- kind & typemask, false,
+ kind & typemask, false, false,
cbufp, refcount_set);
}
else
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-1.c
new file mode 100644
index 0000000..f2e7293
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-1.c
@@ -0,0 +1,31 @@
+#ifdef __cplusplus
+extern "C"
+#else
+extern
+#endif
+void abort (void);
+
+int
+main (void)
+{
+ #define N 5
+ int array[N][N];
+
+ for (int i = 0; i < N; i++)
+ {
+ #pragma omp target enter data map(alloc: array[i:1][0:N])
+
+ #pragma omp target
+ for (int j = 0; j < N; j++)
+ array[i][j] = i + j;
+
+ #pragma omp target exit data map(from: array[i:1][0:N])
+ }
+
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < N; j++)
+ if (array[i][j] != i + j)
+ abort ();
+
+ return 0;
+}
--
2.8.1
^ permalink raw reply [flat|nested] 9+ messages in thread
* Re: [PATCH, OG10, OpenMP 5.0, committed] Implement relaxation of implicit map vs. existing device mappings
2021-05-05 15:17 [PATCH, OG10, OpenMP 5.0, committed] Implement relaxation of implicit map vs. existing device mappings Chung-Lin Tang
@ 2021-05-07 12:35 ` Thomas Schwinge
2021-05-10 9:35 ` Chung-Lin Tang
2021-05-14 13:20 ` [PATCH, OpenMP 5.0] Implement relaxation of implicit map vs. existing device mappings (for mainline trunk) Chung-Lin Tang
0 siblings, 2 replies; 9+ messages in thread
From: Thomas Schwinge @ 2021-05-07 12:35 UTC (permalink / raw)
To: Chung-Lin Tang; +Cc: gcc-patches, Catherine Moore, Tobias Burnus
[-- Attachment #1: Type: text/plain, Size: 54258 bytes --]
Hi Chung-Lin!
On 2021-05-05T23:17:25+0800, Chung-Lin Tang via Gcc-patches <gcc-patches@gcc.gnu.org> wrote:
> This patch implements relaxing the requirements when a map with the implicit attribute encounters
> an overlapping existing map. [...]
Oh, oh, these data mapping interfaces/semantics ares getting more and
more "convoluted"... %-\ (Not your fault, of course.)
Haven't looked in too much detail in the patch/implementation (I'm not
very well-versend in the exact OpenMP semantics anyway), but I suppose we
should do similar things for OpenACC, too. I think we even currently do
have a gimplification-level "hack" to replicate data clauses' array
bounds for implicit data clauses on compute constructs, if the default
"complete" mapping is going to clash with a "limited" mapping that's
specified in an outer OpenACC 'data' directive. (That, of course,
doesn't work for the general case of non-lexical scoping, or dynamic
OpenACC 'enter data', etc., I suppose) I suppose your method could easily
replace and improve that; we shall look into that later.
That said, in your patch, is this current implementation (explicitly)
meant or not meant to be active for OpenACC, too, or just OpenMP (I
couldn't quickly tell), and/or is it (implicitly?) a no-op for OpenACC?
> As the OpenMP 5.0 spec describes on page 320, lines 18-27 (and 5.1 spec,
> page 352, lines 13-22):
>
> "If a single contiguous part of the original storage of a list item with an implicit data-mapping
> attribute has corresponding storage in the device data environment prior to a task encountering the
> construct that is associated with the map clause, only that part of the original storage will have
> corresponding storage in the device data environment as a result of the map clause."
>
> Also tracked in the OpenMP spec context as issue #1463:
> https://github.com/OpenMP/spec/issues/1463
>
> The implementation inside the compiler is to of course, tag the implicitly created maps with some
> indication of "implicit". I've done this with a OMP_CLAUSE_MAP_IMPLICIT_P macro, using
> 'base.deprecated_flag' underneath.
>
> There is an encoding of this as GOMP_MAP_IMPLICIT == GOMP_MAP_FLAG_SPECIAL_3|GOMP_MAP_FLAG_SPECIAL_4
> in include/gomp-constants.h for the runtime, but I've intentionally avoided exploding the entire
> gimplify/omp-low with a new set of GOMP_MAP_IMPLICIT_TO/FROM/etc. symbols, instead adding in the new
> flag bits only at the final runtime call generation during omp-lowering.
>
> The rest is libgomp mapping taking care of the implicit case: allowing map success if an existing
> map is a proper subset of the new map, if the new map is implicit. Straightforward enough I think.
Seems so -- based on my very quick look. ;-)
> There are also some additions to print the implicit attribute during tree pretty-printing, for that
> reason some scan tests were updated.
ACK, thanks.
> Also, another adjustment in this patch is how implicitly created clauses are added to the current
> clause list in gimplify_adjust_omp_clauses(). Instead of simply appending the new clauses to the end,
> this patch adds them at the position "after initial non-map clauses, but right before any existing
> map clauses".
Probably you haven't been testing such a configuration; I've just pushed
"Fix up 'c-c++-common/goacc/firstprivate-mappings-1.c' for C, non-LP64"
to devel/omp/gcc-10 branch in commit
c51cc3b96f0b562deaffcfbcc51043aed216801a, see attached.
> The reason for this is: when combined with other map clauses, for example:
>
> #pragma omp target map(rec.ptr[:N])
> for (int i = 0; i < N; i++)
> rec.ptr[i] += 1;
>
> There will be an implicit map created for map(rec), because of the access inside the target region.
> The expectation is that 'rec' is implicitly mapped, and then the pointed array-section part by 'rec.ptr'
> will be mapped, and then attachment to the 'rec.ptr' field of the mapped 'rec' (in that order).
>
> If the implicit 'map(rec)' is appended to the end, instead of placed before other maps, the attachment
> operation will not find anything to attach to, and the entire region will fail.
But that doesn't (negatively) affect user-visible semantics (OpenMP, and
also OpenACC, if applicable), in that more/bigger objects then get mapped
than were before? (I suppose not?)
Please make sure to put any rationale (like you've posted above) into
source code comments ('gcc/gimplify.c:gimplify_adjust_omp_clauses', I
suppose), or even GCC internals (?) manual, if applicable.
Grüße
Thomas
> Note: this touches a bit on another issue which I will be sending a patch for later:
> per the discussion on omp-lang, an array section list item should *not* be mapping its base-pointer
> (although an attachment attempt should exist), while in current GCC behavior, for struct member pointers
> like 'rec.ptr' above, we do map it (which should be deemed incorrect).
>
> This means that as of right now, this modification of map order doesn't really exhibit the above mentioned
> behavior yet. I have included it as part of this patch because the "[implicit]" tree printing requires
> modifying many gimple scan tests already, so including the test modifications together seems more
> manageable patch-wise.
>
> Tested with no regressions, and pushed to devel/omp/gcc-10. Will be submitting a mainline trunk version later.
>
> Chung-Lin
>
> 2021-05-05 Chung-Lin Tang <cltang@codesourcery.com>
>
> include/ChangeLog:
>
> * gomp-constants.h (GOMP_MAP_IMPLICIT): New special map kind bits value.
> (GOMP_MAP_FLAG_SPECIAL_BITS): Define helper mask for whole set of
> special map kind bits.
> (GOMP_MAP_NONCONTIG_ARRAY_P): Adjust test for non-contiguous array map
> kind bits to be more specific.
> (GOMP_MAP_IMPLICIT_P): New predicate macro for implicit map kinds.
>
> gcc/ChangeLog:
>
> * tree.h (OMP_CLAUSE_MAP_IMPLICIT_P): New access macro for 'implicit'
> bit, using 'base.deprecated_flag' field of tree_node.
> * tree-pretty-print.c (dump_omp_clause): Add support for printing
> implicit attribute in tree dumping.
> * gimplify.c (gimplify_adjust_omp_clauses_1):
> Set OMP_CLAUSE_MAP_IMPLICIT_P to 1 if map clause is implicitly created.
> (gimplify_adjust_omp_clauses): Adjust place of adding implicitly created
> clauses, from simple append, to starting of list, after non-map clauses.
> * omp-low.c (lower_omp_target): Add GOMP_MAP_IMPLICIT bits into kind
> values passed to libgomp for implicit maps.
>
> gcc/testsuite/ChangeLog:
>
> * c-c++-common/gomp/target-implicit-map-1.c: New test.
> * c-c++-common/goacc/combined-reduction.c: Adjust scan test pattern.
> * c-c++-common/goacc/firstprivate-mappings-1.c: Likewise.
> * c-c++-common/goacc/mdc-1.c: Likewise.
> * c-c++-common/goacc/reduction-1.c: Likewise.
> * c-c++-common/goacc/reduction-2.c: Likewise.
> * c-c++-common/goacc/reduction-3.c: Likewise.
> * c-c++-common/goacc/reduction-4.c: Likewise.
> * c-c++-common/goacc/reduction-8.c: Likewise.
> * g++.dg/goacc/firstprivate-mappings-1.C: Likewise.
> * g++.dg/gomp/target-lambda-1.C: Likewise.
> * g++.dg/gomp/target-this-3.C: Likewise.
> * g++.dg/gomp/target-this-4.C: Likewise.
> * gfortran.dg/goacc/common-block-3.f90: Likewise.
> * gfortran.dg/goacc/loop-tree-1.f90: Likewise.
> * gfortran.dg/goacc/private-explicit-kernels-1.f95: Likewise.
> * gfortran.dg/goacc/private-predetermined-kernels-1.f95: Likewise.
>
> libgomp/ChangeLog:
>
> * target.c (gomp_map_vars_existing): Add 'bool implicit' parameter, add
> implicit map handling to allow a "superset" existing map as valid case.
> (get_kind): Adjust to filter out GOMP_MAP_IMPLICIT bits in return value.
> (get_implicit): New function to extract implicit status.
> (gomp_map_fields_existing): Adjust arguments in calls to
> gomp_map_vars_existing, and add uses of get_implicit.
> (gomp_map_vars_internal): Likewise.
>
> * testsuite/libgomp.c-c++-common/target-implicit-map-1.c: New test.
>
>
>
> From a70b5b1aa8b3d32f6728dbfcfc00b0cff8c5219d Mon Sep 17 00:00:00 2001
> From: Chung-Lin Tang <cltang@codesourcery.com>
> Date: Wed, 5 May 2021 08:11:19 -0700
> Subject: [PATCH] OpenMP 5.0: Implement relaxation of implicit map vs. existing
> device mappings
>
> This patch implements relaxing the requirements when a map with the implicit
> attribute encounters an overlapping existing map. As the OpenMP 5.0 spec
> describes on page 320, lines 18-27 (and 5.1 spec, page 352, lines 13-22):
>
> "If a single contiguous part of the original storage of a list item with an
> implicit data-mapping attribute has corresponding storage in the device data
> environment prior to a task encountering the construct that is associated with
> the map clause, only that part of the original storage will have corresponding
> storage in the device data environment as a result of the map clause."
>
> Also tracked in the OpenMP spec context as issue #1463:
> https://github.com/OpenMP/spec/issues/1463
>
> 2021-05-05 Chung-Lin Tang <cltang@codesourcery.com>
>
> include/ChangeLog:
>
> * gomp-constants.h (GOMP_MAP_IMPLICIT): New special map kind bits value.
> (GOMP_MAP_FLAG_SPECIAL_BITS): Define helper mask for whole set of
> special map kind bits.
> (GOMP_MAP_NONCONTIG_ARRAY_P): Adjust test for non-contiguous array map
> kind bits to be more specific.
> (GOMP_MAP_IMPLICIT_P): New predicate macro for implicit map kinds.
>
> gcc/ChangeLog:
>
> * tree.h (OMP_CLAUSE_MAP_IMPLICIT_P): New access macro for 'implicit'
> bit, using 'base.deprecated_flag' field of tree_node.
> * tree-pretty-print.c (dump_omp_clause): Add support for printing
> implicit attribute in tree dumping.
> * gimplify.c (gimplify_adjust_omp_clauses_1):
> Set OMP_CLAUSE_MAP_IMPLICIT_P to 1 if map clause is implicitly created.
> (gimplify_adjust_omp_clauses): Adjust place of adding implicitly created
> clauses, from simple append, to starting of list, after non-map clauses.
> * omp-low.c (lower_omp_target): Add GOMP_MAP_IMPLICIT bits into kind
> values passed to libgomp for implicit maps.
>
> gcc/testsuite/ChangeLog:
>
> * c-c++-common/gomp/target-implicit-map-1.c: New test.
> * c-c++-common/goacc/combined-reduction.c: Adjust scan test pattern.
> * c-c++-common/goacc/firstprivate-mappings-1.c: Likewise.
> * c-c++-common/goacc/mdc-1.c: Likewise.
> * c-c++-common/goacc/reduction-1.c: Likewise.
> * c-c++-common/goacc/reduction-2.c: Likewise.
> * c-c++-common/goacc/reduction-3.c: Likewise.
> * c-c++-common/goacc/reduction-4.c: Likewise.
> * c-c++-common/goacc/reduction-8.c: Likewise.
> * g++.dg/goacc/firstprivate-mappings-1.C: Likewise.
> * g++.dg/gomp/target-lambda-1.C: Likewise.
> * g++.dg/gomp/target-this-3.C: Likewise.
> * g++.dg/gomp/target-this-4.C: Likewise.
> * gfortran.dg/goacc/common-block-3.f90: Likewise.
> * gfortran.dg/goacc/loop-tree-1.f90: Likewise.
> * gfortran.dg/goacc/private-explicit-kernels-1.f95: Likewise.
> * gfortran.dg/goacc/private-predetermined-kernels-1.f95: Likewise.
>
> libgomp/ChangeLog:
>
> * target.c (gomp_map_vars_existing): Add 'bool implicit' parameter, add
> implicit map handling to allow a "superset" existing map as valid case.
> (get_kind): Adjust to filter out GOMP_MAP_IMPLICIT bits in return value.
> (get_implicit): New function to extract implicit status.
> (gomp_map_fields_existing): Adjust arguments in calls to
> gomp_map_vars_existing, and add uses of get_implicit.
> (gomp_map_vars_internal): Likewise.
>
> * testsuite/libgomp.c-c++-common/target-implicit-map-1.c: New test.
> ---
> gcc/gimplify.c | 11 ++-
> gcc/omp-low.c | 13 ++++
> .../c-c++-common/goacc/combined-reduction.c | 4 +-
> .../c-c++-common/goacc/firstprivate-mappings-1.c | 6 +-
> gcc/testsuite/c-c++-common/goacc/mdc-1.c | 2 +-
> gcc/testsuite/c-c++-common/goacc/reduction-1.c | 4 +-
> gcc/testsuite/c-c++-common/goacc/reduction-2.c | 4 +-
> gcc/testsuite/c-c++-common/goacc/reduction-3.c | 4 +-
> gcc/testsuite/c-c++-common/goacc/reduction-4.c | 4 +-
> gcc/testsuite/c-c++-common/goacc/reduction-8.c | 10 +--
> .../c-c++-common/gomp/target-implicit-map-1.c | 39 +++++++++++
> .../g++.dg/goacc/firstprivate-mappings-1.C | 2 +-
> gcc/testsuite/g++.dg/gomp/target-lambda-1.C | 6 +-
> gcc/testsuite/g++.dg/gomp/target-this-3.C | 4 +-
> gcc/testsuite/g++.dg/gomp/target-this-4.C | 4 +-
> gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 | 8 +--
> gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90 | 2 +-
> .../goacc/private-explicit-kernels-1.f95 | 4 +-
> .../goacc/private-predetermined-kernels-1.f95 | 4 +-
> gcc/tree-pretty-print.c | 3 +
> gcc/tree.h | 4 ++
> include/gomp-constants.h | 17 ++++-
> libgomp/target.c | 78 ++++++++++++++++------
> .../libgomp.c-c++-common/target-implicit-map-1.c | 31 +++++++++
> 24 files changed, 211 insertions(+), 57 deletions(-)
> create mode 100644 gcc/testsuite/c-c++-common/gomp/target-implicit-map-1.c
> create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-1.c
>
> diff --git a/gcc/gimplify.c b/gcc/gimplify.c
> index 91aa15d..ba071e8 100644
> --- a/gcc/gimplify.c
> +++ b/gcc/gimplify.c
> @@ -10579,6 +10579,7 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
> gcc_unreachable ();
> }
> OMP_CLAUSE_SET_MAP_KIND (clause, kind);
> + OMP_CLAUSE_MAP_IMPLICIT_P (clause) = 1;
> if (DECL_SIZE (decl)
> && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
> {
> @@ -11158,9 +11159,15 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
> list_p = &OMP_CLAUSE_CHAIN (c);
> }
>
> - /* Add in any implicit data sharing. */
> + /* Add in any implicit data sharing. Implicit clauses are added at the start
> + of the clause list, but after any non-map clauses. */
> struct gimplify_adjust_omp_clauses_data data;
> - data.list_p = list_p;
> + tree *implicit_add_list_p = orig_list_p;
> + while (*implicit_add_list_p
> + && OMP_CLAUSE_CODE (*implicit_add_list_p) != OMP_CLAUSE_MAP)
> + implicit_add_list_p = &OMP_CLAUSE_CHAIN (*implicit_add_list_p);
> +
> + data.list_p = implicit_add_list_p;
> data.pre_p = pre_p;
> splay_tree_foreach (ctx->variables, gimplify_adjust_omp_clauses_1, &data);
>
> diff --git a/gcc/omp-low.c b/gcc/omp-low.c
> index 66519ad..64b7c19 100644
> --- a/gcc/omp-low.c
> +++ b/gcc/omp-low.c
> @@ -12920,6 +12920,19 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
> else if (integer_nonzerop (s))
> tkind_zero = tkind;
> }
> + if (tkind_zero == tkind
> + && OMP_CLAUSE_MAP_IMPLICIT_P (c)
> + && (((tkind & GOMP_MAP_FLAG_SPECIAL_BITS)
> + & ~GOMP_MAP_IMPLICIT)
> + == 0))
> + {
> + /* If this is an implicit map, and the GOMP_MAP_IMPLICIT
> + bits are not interfered by other special bit encodings,
> + then turn the GOMP_IMPLICIT_BIT flag on for the runtime
> + to see. */
> + tkind |= GOMP_MAP_IMPLICIT;
> + tkind_zero = tkind;
> + }
> break;
> case OMP_CLAUSE_FIRSTPRIVATE:
> gcc_checking_assert (is_gimple_omp_oacc (ctx->stmt));
> diff --git a/gcc/testsuite/c-c++-common/goacc/combined-reduction.c b/gcc/testsuite/c-c++-common/goacc/combined-reduction.c
> index ecf23f5..fa67e08 100644
> --- a/gcc/testsuite/c-c++-common/goacc/combined-reduction.c
> +++ b/gcc/testsuite/c-c++-common/goacc/combined-reduction.c
> @@ -23,7 +23,7 @@ main ()
> return 0;
> }
>
> -/* { dg-final { scan-tree-dump-times "omp target oacc_parallel reduction.+:v1. map.tofrom:v1" 1 "gimple" } } */
> +/* { dg-final { scan-tree-dump-times "omp target oacc_parallel reduction.+:v1. firstprivate.n. map.tofrom:v1" 1 "gimple" } } */
> /* { dg-final { scan-tree-dump-times "acc loop reduction.+:v1. private.i." 1 "gimple" } } */
> -/* { dg-final { scan-tree-dump-times "omp target oacc_kernels map.force_tofrom:n .len: 4.. map.force_tofrom:v1 .len: 4.." 1 "gimple" } } */
> +/* { dg-final { scan-tree-dump-times "omp target oacc_kernels map.force_tofrom:n .len: 4..implicit.. map.force_tofrom:v1 .len: 4..implicit.." 1 "gimple" } } */
> /* { dg-final { scan-tree-dump-times "acc loop reduction.+:v1. private.i." 1 "gimple" } } */
> diff --git a/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c b/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c
> index 7987bea..f43e4b4 100644
> --- a/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c
> +++ b/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c
> @@ -419,12 +419,12 @@ vla (int array_li)
> copyout (array_so)
> /* The gimplifier has created an implicit 'firstprivate' clause for the array
> length.
> - { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel map\(from:array_so \[len: 4\]\) firstprivate\(array_li.[0-9]+\)} omplower { target { ! c++ } } } }
> - { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel map\(from:array_so \[len: 4\]\) firstprivate\(} omplower { target { c++ } } } }
> + { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel firstprivate\(array_li.[0-9]+\) map\(from:array_so \[len: 4\]\)} omplower { target { ! c++ } } } }
> + { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel firstprivate\([^)]+\) map\(from:array_so \[len: 4\]\)} omplower { target { c++ } } } }
> (C++ computes an intermediate value, so can't scan for 'firstprivate(array_li)'.) */
> /* For C, non-LP64, the gimplifier has also created a mapping for the array
> itself; PR90859.
> - { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel map\(from:array_so \[len: 4\]\) firstprivate\(array_li.[0-9]+\) map\(tofrom:\(\*array.[0-9]+\) \[len: D\.[0-9]+\]\) map\(firstprivate:array \[pointer assign, bias: 0\]\) \[} omplower { target { c && { ! lp64 } } } } } */
> + { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel firstprivate\(array_li.[0-9]+\) map\(tofrom:\(\*array.[0-9]+\) \[len: D\.[0-9]+\]\[implicit\]\) map\(firstprivate:array \[pointer assign, bias: 0\]\) map\(from:array_so \[len: 4\]\) \[} omplower { target { c && { ! lp64 } } } } } */
> {
> array_so = sizeof array;
> }
> diff --git a/gcc/testsuite/c-c++-common/goacc/mdc-1.c b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
> index 337c1f7..9f43de4 100644
> --- a/gcc/testsuite/c-c++-common/goacc/mdc-1.c
> +++ b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
> @@ -45,7 +45,7 @@ t1 ()
>
> /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:s .len: 32.." 1 "omplower" } } */
> /* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.tofrom:.z .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len: 8.. map.tofrom:._1 .len: 40.. map.attach:s.a .bias: 0.." 1 "omplower" } } */
> -/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.attach:s.e .bias: 0.. map.tofrom:s .len: 32" 1 "omplower" } } */
> +/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.tofrom:s .len: 32..implicit.. map.attach:s.e .bias: 0.." 1 "omplower" } } */
> /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.attach:a .bias: 0.." 1 "omplower" } } */
> /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:a .bias: 0.." 1 "omplower" } } */
> /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:a .len: 8.." 1 "omplower" } } */
> diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-1.c b/gcc/testsuite/c-c++-common/goacc/reduction-1.c
> index 35bfc86..d9e3c38 100644
> --- a/gcc/testsuite/c-c++-common/goacc/reduction-1.c
> +++ b/gcc/testsuite/c-c++-common/goacc/reduction-1.c
> @@ -68,5 +68,5 @@ main(void)
> }
>
> /* Check that default copy maps are generated for loop reductions. */
> -/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\]\\)" 7 "gimple" } } */
> -/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
> +/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 7 "gimple" } } */
> +/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 2 "gimple" } } */
> diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-2.c b/gcc/testsuite/c-c++-common/goacc/reduction-2.c
> index 9dba035..18dc03c 100644
> --- a/gcc/testsuite/c-c++-common/goacc/reduction-2.c
> +++ b/gcc/testsuite/c-c++-common/goacc/reduction-2.c
> @@ -50,5 +50,5 @@ main(void)
> }
>
> /* Check that default copy maps are generated for loop reductions. */
> -/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\]\\)" 4 "gimple" } } */
> -/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
> +/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 4 "gimple" } } */
> +/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 2 "gimple" } } */
> diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-3.c b/gcc/testsuite/c-c++-common/goacc/reduction-3.c
> index 669cd43..2311d4b 100644
> --- a/gcc/testsuite/c-c++-common/goacc/reduction-3.c
> +++ b/gcc/testsuite/c-c++-common/goacc/reduction-3.c
> @@ -50,5 +50,5 @@ main(void)
> }
>
> /* Check that default copy maps are generated for loop reductions. */
> -/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\]\\)" 4 "gimple" } } */
> -/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
> +/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 4 "gimple" } } */
> +/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 2 "gimple" } } */
> diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-4.c b/gcc/testsuite/c-c++-common/goacc/reduction-4.c
> index 5c3dfb1..57823f8 100644
> --- a/gcc/testsuite/c-c++-common/goacc/reduction-4.c
> +++ b/gcc/testsuite/c-c++-common/goacc/reduction-4.c
> @@ -38,5 +38,5 @@ main(void)
> }
>
> /* Check that default copy maps are generated for loop reductions. */
> -/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
> -/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
> +/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 2 "gimple" } } */
> +/* { dg-final { scan-tree-dump-times "map\\(tofrom:lresult \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 2 "gimple" } } */
> diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-8.c b/gcc/testsuite/c-c++-common/goacc/reduction-8.c
> index 8a0283f..8494e59 100644
> --- a/gcc/testsuite/c-c++-common/goacc/reduction-8.c
> +++ b/gcc/testsuite/c-c++-common/goacc/reduction-8.c
> @@ -87,8 +87,10 @@ main(void)
>
> /* Check that default copy maps are generated for loop reductions. */
> /* { dg-final { scan-tree-dump-times "reduction..:result. map.tofrom:result .len: 4.." 1 "gimple" } } */
> -/* { dg-final { scan-tree-dump-times "oacc_parallel map.tofrom:result .len: 4.." 2 "gimple" } } */
> -/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000.. firstprivate.result." 3 "gimple" } } */
> -/* { dg-final { scan-tree-dump-times "map.tofrom:result .len: 4.. map.tofrom:array .len: 4000.." 1 "gimple" } } */
> -/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000.. map.force_tofrom:result .len: 4.." 1 "gimple" } } */
> +
> +/* { dg-final { scan-tree-dump-times "oacc_parallel map.tofrom:result .len: 4..implicit.." 1 "gimple" } } */
> +
> +/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000..implicit.. firstprivate.result." 3 "gimple" } } */
> +/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000..implicit.. map.tofrom:result .len: 4.." 1 "gimple" } } */
> +/* { dg-final { scan-tree-dump-times "map.tofrom:array .len: 4000..implicit.. map.force_tofrom:result .len: 4..implicit.." 1 "gimple" } } */
>
> diff --git a/gcc/testsuite/c-c++-common/gomp/target-implicit-map-1.c b/gcc/testsuite/c-c++-common/gomp/target-implicit-map-1.c
> new file mode 100644
> index 0000000..52944fd
> --- /dev/null
> +++ b/gcc/testsuite/c-c++-common/gomp/target-implicit-map-1.c
> @@ -0,0 +1,39 @@
> +/* { dg-do compile } */
> +/* { dg-additional-options "-fdump-tree-gimple" } */
> +#ifdef __cplusplus
> +extern "C"
> +#else
> +extern
> +#endif
> +void abort (void);
> +
> +int
> +main (void)
> +{
> + #define N 5
> + int array[N][N];
> +
> + for (int i = 0; i < N; i++)
> + {
> + #pragma omp target enter data map(alloc: array[i:1][0:N])
> +
> + #pragma omp target
> + for (int j = 0; j < N; j++)
> + array[i][j] = i * 10 + j;
> +
> + #pragma omp target exit data map(from: array[i:1][0:N])
> + }
> +
> + for (int i = 0; i < N; i++)
> + for (int j = 0; j < N; j++)
> + if (array[i][j] != i + j)
> + abort ();
> +
> + return 0;
> +}
> +
> +/* { dg-final { scan-tree-dump {#pragma omp target enter data map\(alloc:array\[[^]]+\]\[0\] \[len: [0-9]+\]\)} "gimple" } } */
> +
> +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(i\) map\(tofrom:array \[len: [0-9]+\]\[implicit\]\)} "gimple" } } */
> +
> +/* { dg-final { scan-tree-dump {#pragma omp target exit data map\(from:array\[[^]]+\]\[0\] \[len: [0-9]+\]\)} "gimple" } } */
> diff --git a/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C b/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C
> index 1b1badb..99a3bd4 100644
> --- a/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C
> +++ b/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C
> @@ -416,7 +416,7 @@ vla (int &array_li)
> copyout (array_so)
> /* The gimplifier has created an implicit 'firstprivate' clause for the array
> length.
> - { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel map\(from:array_so \[len: 4\]\) firstprivate\(} omplower } }
> + { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel firstprivate\([^)]+\) map\(from:array_so \[len: 4\]\)} omplower } }
> (C++ computes an intermediate value, so can't scan for 'firstprivate(array_li)'.) */
> {
> array_so = sizeof array;
> diff --git a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
> index 7dceef8..e5a24d7 100644
> --- a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
> +++ b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
> @@ -87,8 +87,8 @@ int main (void)
> return 0;
> }
>
> -/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */
> +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */
>
> -/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(end\) firstprivate\(begin\)} "gimple" } } */
> +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */
>
> -/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data2 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(end\) firstprivate\(begin\)} "gimple" } } */
> +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data2 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */
> diff --git a/gcc/testsuite/g++.dg/gomp/target-this-3.C b/gcc/testsuite/g++.dg/gomp/target-this-3.C
> index 08568f9..2755b4b 100644
> --- a/gcc/testsuite/g++.dg/gomp/target-this-3.C
> +++ b/gcc/testsuite/g++.dg/gomp/target-this-3.C
> @@ -100,6 +100,6 @@ int main (void)
> return 0;
> }
>
> -/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:this->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9+] \[len: 0\]\) firstprivate\(n\)} "gimple" } } */
> +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:this->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9+] \[len: 0\]\)} "gimple" } } */
>
> -/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:this->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(n\)} "gimple" } } */
> +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:this->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */
> diff --git a/gcc/testsuite/g++.dg/gomp/target-this-4.C b/gcc/testsuite/g++.dg/gomp/target-this-4.C
> index 3b2d581..3703762 100644
> --- a/gcc/testsuite/g++.dg/gomp/target-this-4.C
> +++ b/gcc/testsuite/g++.dg/gomp/target-this-4.C
> @@ -102,6 +102,6 @@ int main (void)
> return 0;
> }
>
> -/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: 1\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */
> +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: 1\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */
>
> -/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:_[0-9]+->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */
> +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:_[0-9]+->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */
> diff --git a/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 b/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90
> index e43d376..e9f169f 100644
> --- a/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90
> +++ b/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90
> @@ -33,10 +33,10 @@ end program main
>
> ! { dg-final { scan-tree-dump-times "omp target oacc_data_kernels .*map\\(tofrom:x \\\[len: 400\\\]\\)" 1 "omplower" } }
> ! { dg-final { scan-tree-dump-times "omp target oacc_data_kernels .*map\\(tofrom:y \\\[len: 400\\\]\\\)" 1 "omplower" } }
> -! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:x \\\[len: 400\\\]\\)" 1 "omplower" } }
> -! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:y \\\[len: 400\\\]\\\)" 1 "omplower" } }
> -! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_tofrom:i \\\[len: 4\\\]\\)" 1 "omplower" } }
> -! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_tofrom:c \\\[len: 4\\\]\\)" 1 "omplower" } }
> +! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:x \\\[len: 400\\\]\\\[implicit\\\]\\)" 1 "omplower" } }
> +! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:y \\\[len: 400\\\]\\\[implicit\\\]\\\)" 1 "omplower" } }
> +! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_tofrom:i \\\[len: 4\\\]\\\[implicit\\\]\\)" 1 "omplower" } }
> +! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_tofrom:c \\\[len: 4\\\]\\\[implicit\\\]\\)" 1 "omplower" } }
>
> ! Expecting no mapping of un-referenced common-blocks variables
>
> diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90 b/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90
> index 150f930..4cdfc55 100644
> --- a/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90
> +++ b/gcc/testsuite/gfortran.dg/goacc/loop-tree-1.f90
> @@ -44,4 +44,4 @@ end program test
>
> ! { dg-final { scan-tree-dump-times "private\\(m\\)" 1 "original" } }
> ! { dg-final { scan-tree-dump-times "reduction\\(\\+:sum\\)" 1 "original" } }
> -! { dg-final { scan-tree-dump-times "map\\(tofrom:sum \\\[len: \[0-9\]+\\\]\\)" 1 "gimple" } }
> +! { dg-final { scan-tree-dump-times "map\\(tofrom:sum \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 1 "gimple" } }
> diff --git a/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95 b/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95
> index 0c47045..fef5126 100644
> --- a/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95
> +++ b/gcc/testsuite/gfortran.dg/goacc/private-explicit-kernels-1.f95
> @@ -83,7 +83,7 @@ program test
> !$acc kernels ! Explicit "private(i2_2_s)" clause cannot be specified here.
> ! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "original" { xfail *-*-* } } } ! PR90067
> ! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
> - ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i2_2_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
> + ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i2_2_s \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
> do i2_2_s = 1, 100
> !$acc loop private(j2_2_s) independent
> ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_2_s\\) independent" 1 "original" } }
> @@ -234,7 +234,7 @@ program test
> !$acc kernels ! Explicit "private(i3_5_s)" clause cannot be specified here.
> ! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "original" { xfail *-*-* } } } ! PR90067
> ! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
> - ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i3_5_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
> + ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i3_5_s \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
> do i3_5_s = 1, 100
> !$acc loop private(j3_5_s) independent
> ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_5_s\\) independent" 1 "original" } }
> diff --git a/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95 b/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95
> index 3357a20..38459cf 100644
> --- a/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95
> +++ b/gcc/testsuite/gfortran.dg/goacc/private-predetermined-kernels-1.f95
> @@ -83,7 +83,7 @@ program test
> !$acc kernels
> ! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "original" { xfail *-*-* } } } ! PR90067
> ! { dg-final { scan-tree-dump-times "private\\(i2_2_s\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
> - ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i2_2_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
> + ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i2_2_s \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
> do i2_2_s = 1, 100
> !$acc loop independent
> ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j2_2_s\\) independent" 1 "original" } }
> @@ -234,7 +234,7 @@ program test
> !$acc kernels
> ! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "original" { xfail *-*-* } } } ! PR90067
> ! { dg-final { scan-tree-dump-times "private\\(i3_5_s\\)" 1 "gimple" { xfail *-*-* } } } ! PR90067
> - ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i3_5_s \\\[len: \[0-9\]+\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
> + ! { dg-final { scan-tree-dump-times "#pragma omp target oacc_kernels map\\(force_tofrom:i3_5_s \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" 0 "gimple" { xfail *-*-* } } } ! PR90067
> do i3_5_s = 1, 100
> !$acc loop independent
> ! { dg-final { scan-tree-dump-times "#pragma acc loop private\\(j3_5_s\\) independent" 1 "original" } }
> diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
> index 261cc9d..4cd4606 100644
> --- a/gcc/tree-pretty-print.c
> +++ b/gcc/tree-pretty-print.c
> @@ -946,6 +946,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
> spc, flags, false);
> pp_right_bracket (pp);
> }
> + if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
> + && OMP_CLAUSE_MAP_IMPLICIT_P (clause))
> + pp_string (pp, "[implicit]");
> pp_right_paren (pp);
> break;
>
> diff --git a/gcc/tree.h b/gcc/tree.h
> index 8d9829c..647c5ba 100644
> --- a/gcc/tree.h
> +++ b/gcc/tree.h
> @@ -1637,6 +1637,10 @@ class auto_suppress_location_wrappers
> variable. */
> #define OMP_CLAUSE_MAP_IN_REDUCTION(NODE) \
> TREE_PRIVATE (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
> +/* Nonzero if this map clause was created through implicit data-mapping
> + rules. */
> +#define OMP_CLAUSE_MAP_IMPLICIT_P(NODE) \
> + (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.deprecated_flag)
>
> /* True on an OMP_CLAUSE_USE_DEVICE_PTR with an OpenACC 'if_present'
> clause. */
> diff --git a/include/gomp-constants.h b/include/gomp-constants.h
> index b8efb30..33cfcb9 100644
> --- a/include/gomp-constants.h
> +++ b/include/gomp-constants.h
> @@ -46,6 +46,16 @@
> | GOMP_MAP_FLAG_SPECIAL_0)
> #define GOMP_MAP_DEEP_COPY (GOMP_MAP_FLAG_SPECIAL_4 \
> | GOMP_MAP_FLAG_SPECIAL_2)
> +/* This value indicates the map was created implicitly according to
> + OpenMP rules. */
> +#define GOMP_MAP_IMPLICIT (GOMP_MAP_FLAG_SPECIAL_3 \
> + | GOMP_MAP_FLAG_SPECIAL_4)
> +/* Mask for entire set of special map kind bits. */
> +#define GOMP_MAP_FLAG_SPECIAL_BITS (GOMP_MAP_FLAG_SPECIAL_0 \
> + | GOMP_MAP_FLAG_SPECIAL_1 \
> + | GOMP_MAP_FLAG_SPECIAL_2 \
> + | GOMP_MAP_FLAG_SPECIAL_3 \
> + | GOMP_MAP_FLAG_SPECIAL_4)
> /* Flag to force a specific behavior (or else, trigger a run-time error). */
> #define GOMP_MAP_FLAG_FORCE (1 << 7)
>
> @@ -225,7 +235,12 @@ enum gomp_map_kind
> (GOMP_MAP_ALWAYS_TO_P (X) || ((X) == GOMP_MAP_ALWAYS_FROM))
>
> #define GOMP_MAP_NONCONTIG_ARRAY_P(X) \
> - ((X) & GOMP_MAP_NONCONTIG_ARRAY)
> + (((X) & GOMP_MAP_FLAG_SPECIAL_BITS) == GOMP_MAP_NONCONTIG_ARRAY \
> + || (X) == GOMP_MAP_NONCONTIG_ARRAY_FORCE_PRESENT)
> +
> +#define GOMP_MAP_IMPLICIT_P(X) \
> + (((X) & GOMP_MAP_FLAG_SPECIAL_BITS) == GOMP_MAP_IMPLICIT)
> +
>
> /* Asynchronous behavior. Keep in sync with
> libgomp/{openacc.h,openacc.f90,openacc_lib.h}:acc_async_t. */
> diff --git a/libgomp/target.c b/libgomp/target.c
> index 9c75826..ecda2ef 100644
> --- a/libgomp/target.c
> +++ b/libgomp/target.c
> @@ -510,7 +510,7 @@ static inline void
> gomp_map_vars_existing (struct gomp_device_descr *devicep,
> struct goacc_asyncqueue *aq, splay_tree_key oldn,
> splay_tree_key newn, struct target_var_desc *tgt_var,
> - unsigned char kind, bool always_to_flag,
> + unsigned char kind, bool always_to_flag, bool implicit,
> struct gomp_coalesce_buf *cbuf,
> htab_t *refcount_set)
> {
> @@ -522,11 +522,22 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep,
> tgt_var->always_copy_from = GOMP_MAP_ALWAYS_FROM_P (kind);
> tgt_var->is_attach = false;
> tgt_var->offset = newn->host_start - oldn->host_start;
> - tgt_var->length = newn->host_end - newn->host_start;
> +
> + /* For implicit maps, old contained in new is valid. */
> + bool implicit_subset = (implicit
> + && newn->host_start <= oldn->host_start
> + && oldn->host_end <= newn->host_end);
> + if (implicit_subset)
> + tgt_var->length = oldn->host_end - oldn->host_start;
> + else
> + tgt_var->length = newn->host_end - newn->host_start;
>
> if ((kind & GOMP_MAP_FLAG_FORCE)
> - || oldn->host_start > newn->host_start
> - || oldn->host_end < newn->host_end)
> + /* For implicit maps, old contained in new is valid. */
> + || !(implicit_subset
> + /* Otherwise, new contained inside old is considered valid. */
> + || (oldn->host_start <= newn->host_start
> + && newn->host_end <= oldn->host_end)))
> {
> gomp_mutex_unlock (&devicep->lock);
> gomp_fatal ("Trying to map into device [%p..%p) object when "
> @@ -536,11 +547,17 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep,
> }
>
> if (GOMP_MAP_ALWAYS_TO_P (kind) || always_to_flag)
> - gomp_copy_host2dev (devicep, aq,
> - (void *) (oldn->tgt->tgt_start + oldn->tgt_offset
> - + newn->host_start - oldn->host_start),
> - (void *) newn->host_start,
> - newn->host_end - newn->host_start, false, cbuf);
> + {
> + /* Implicit + always should not happen. If this does occur, below
> + address/length adjustment is a TODO. */
> + assert (!implicit_subset);
> +
> + gomp_copy_host2dev (devicep, aq,
> + (void *) (oldn->tgt->tgt_start + oldn->tgt_offset
> + + newn->host_start - oldn->host_start),
> + (void *) newn->host_start,
> + newn->host_end - newn->host_start, false, cbuf);
> + }
>
> gomp_increment_refcount (oldn, refcount_set);
> }
> @@ -548,8 +565,24 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep,
> static int
> get_kind (bool short_mapkind, void *kinds, int idx)
> {
> - return short_mapkind ? ((unsigned short *) kinds)[idx]
> - : ((unsigned char *) kinds)[idx];
> + int val = (short_mapkind
> + ? ((unsigned short *) kinds)[idx]
> + : ((unsigned char *) kinds)[idx]);
> +
> + if (GOMP_MAP_IMPLICIT_P (val))
> + val &= ~GOMP_MAP_IMPLICIT;
> + return val;
> +}
> +
> +
> +static bool
> +get_implicit (bool short_mapkind, void *kinds, int idx)
> +{
> + int val = (short_mapkind
> + ? ((unsigned short *) kinds)[idx]
> + : ((unsigned char *) kinds)[idx]);
> +
> + return GOMP_MAP_IMPLICIT_P (val);
> }
>
> static void
> @@ -612,6 +645,7 @@ gomp_map_fields_existing (struct target_mem_desc *tgt,
> struct splay_tree_s *mem_map = &devicep->mem_map;
> struct splay_tree_key_s cur_node;
> int kind;
> + bool implicit;
> const bool short_mapkind = true;
> const int typemask = short_mapkind ? 0xff : 0x7;
>
> @@ -619,12 +653,14 @@ gomp_map_fields_existing (struct target_mem_desc *tgt,
> cur_node.host_end = cur_node.host_start + sizes[i];
> splay_tree_key n2 = splay_tree_lookup (mem_map, &cur_node);
> kind = get_kind (short_mapkind, kinds, i);
> + implicit = get_implicit (short_mapkind, kinds, i);
> if (n2
> && n2->tgt == n->tgt
> && n2->host_start - n->host_start == n2->tgt_offset - n->tgt_offset)
> {
> gomp_map_vars_existing (devicep, aq, n2, &cur_node, &tgt->list[i],
> - kind & typemask, false, cbuf, refcount_set);
> + kind & typemask, false, implicit, cbuf,
> + refcount_set);
> return;
> }
> if (sizes[i] == 0)
> @@ -640,7 +676,8 @@ gomp_map_fields_existing (struct target_mem_desc *tgt,
> == n2->tgt_offset - n->tgt_offset)
> {
> gomp_map_vars_existing (devicep, aq, n2, &cur_node, &tgt->list[i],
> - kind & typemask, false, cbuf, refcount_set);
> + kind & typemask, false, implicit, cbuf,
> + refcount_set);
> return;
> }
> }
> @@ -652,7 +689,8 @@ gomp_map_fields_existing (struct target_mem_desc *tgt,
> && n2->host_start - n->host_start == n2->tgt_offset - n->tgt_offset)
> {
> gomp_map_vars_existing (devicep, aq, n2, &cur_node, &tgt->list[i],
> - kind & typemask, false, cbuf, refcount_set);
> + kind & typemask, false, implicit, cbuf,
> + refcount_set);
> return;
> }
> }
> @@ -898,6 +936,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
> for (i = 0; i < mapnum; i++)
> {
> int kind = get_kind (short_mapkind, kinds, i);
> + bool implicit = get_implicit (short_mapkind, kinds, i);
> if (hostaddrs[i] == NULL
> || (kind & typemask) == GOMP_MAP_FIRSTPRIVATE_INT)
> {
> @@ -1104,8 +1143,8 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
> }
> }
> gomp_map_vars_existing (devicep, aq, n, &cur_node, &tgt->list[i],
> - kind & typemask, always_to_cnt > 0, NULL,
> - refcount_set);
> + kind & typemask, always_to_cnt > 0, implicit,
> + NULL, refcount_set);
> i += always_to_cnt;
> }
> else
> @@ -1182,7 +1221,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
> {
> assert (n->refcount != REFCOUNT_LINK);
> gomp_map_vars_existing (devicep, aq, n, &cur_node, row_desc,
> - kind & typemask, false,
> + kind & typemask, false, false,
> /* TODO: cbuf? */ NULL, refcount_set);
> }
> else
> @@ -1312,6 +1351,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
> else if (tgt->list[i].key == NULL)
> {
> int kind = get_kind (short_mapkind, kinds, i);
> + bool implicit = get_implicit (short_mapkind, kinds, i);
> if (hostaddrs[i] == NULL)
> continue;
> switch (kind & typemask)
> @@ -1483,7 +1523,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
> splay_tree_key n = splay_tree_lookup (mem_map, k);
> if (n && n->refcount != REFCOUNT_LINK)
> gomp_map_vars_existing (devicep, aq, n, k, &tgt->list[i],
> - kind & typemask, false, cbufp,
> + kind & typemask, false, implicit, cbufp,
> refcount_set);
> else
> {
> @@ -1702,7 +1742,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
> {
> assert (k->refcount != REFCOUNT_LINK);
> gomp_map_vars_existing (devicep, aq, k, &cur_node, row_desc,
> - kind & typemask, false,
> + kind & typemask, false, false,
> cbufp, refcount_set);
> }
> else
> diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-1.c
> new file mode 100644
> index 0000000..f2e7293
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-1.c
> @@ -0,0 +1,31 @@
> +#ifdef __cplusplus
> +extern "C"
> +#else
> +extern
> +#endif
> +void abort (void);
> +
> +int
> +main (void)
> +{
> + #define N 5
> + int array[N][N];
> +
> + for (int i = 0; i < N; i++)
> + {
> + #pragma omp target enter data map(alloc: array[i:1][0:N])
> +
> + #pragma omp target
> + for (int j = 0; j < N; j++)
> + array[i][j] = i + j;
> +
> + #pragma omp target exit data map(from: array[i:1][0:N])
> + }
> +
> + for (int i = 0; i < N; i++)
> + for (int j = 0; j < N; j++)
> + if (array[i][j] != i + j)
> + abort ();
> +
> + return 0;
> +}
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-Fix-up-c-c-common-goacc-firstprivate-mappings-1.og10.patch --]
[-- Type: text/x-diff, Size: 2434 bytes --]
From c51cc3b96f0b562deaffcfbcc51043aed216801a Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Fri, 7 May 2021 10:57:36 +0200
Subject: [PATCH] Fix up 'c-c++-common/goacc/firstprivate-mappings-1.c' for C,
non-LP64
Follow-up to recent og10 commit a70b5b1aa8b3d32f6728dbfcfc00b0cff8c5219d
"OpenMP 5.0: Implement relaxation of implicit map vs. existing device
mappings".
gcc/testsuite/
* c-c++-common/goacc/firstprivate-mappings-1.c: Fix up for C,
non-LP64.
---
gcc/testsuite/ChangeLog.omp | 5 +++++
gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c | 4 ++--
2 files changed, 7 insertions(+), 2 deletions(-)
diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index 45277a37fd2..b18615a1603 100644
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,8 @@
+2021-05-07 Thomas Schwinge <thomas@codesourcery.com>
+
+ * c-c++-common/goacc/firstprivate-mappings-1.c: Fix up for C,
+ non-LP64.
+
2021-04-11 Hafiz Abid Qadeer <abidh@codesourcery.com>
Backport from mainline
diff --git a/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c b/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c
index f43e4b46cb6..ab09dee3d37 100644
--- a/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c
@@ -419,8 +419,8 @@ vla (int array_li)
copyout (array_so)
/* The gimplifier has created an implicit 'firstprivate' clause for the array
length.
- { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel firstprivate\(array_li.[0-9]+\) map\(from:array_so \[len: 4\]\)} omplower { target { ! c++ } } } }
- { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel firstprivate\([^)]+\) map\(from:array_so \[len: 4\]\)} omplower { target { c++ } } } }
+ { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel firstprivate\(array_li.[0-9]+\) map\(from:array_so \[len: 4\]\) \[} omplower { target { c && lp64 } } } }
+ { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel firstprivate\(D\.[0-9]+\) map\(from:array_so \[len: 4\]\) \[} omplower { target { c++ } } } }
(C++ computes an intermediate value, so can't scan for 'firstprivate(array_li)'.) */
/* For C, non-LP64, the gimplifier has also created a mapping for the array
itself; PR90859.
--
2.30.2
^ permalink raw reply [flat|nested] 9+ messages in thread
* Re: [PATCH, OG10, OpenMP 5.0, committed] Implement relaxation of implicit map vs. existing device mappings
2021-05-07 12:35 ` Thomas Schwinge
@ 2021-05-10 9:35 ` Chung-Lin Tang
2021-05-14 13:20 ` [PATCH, OpenMP 5.0] Implement relaxation of implicit map vs. existing device mappings (for mainline trunk) Chung-Lin Tang
1 sibling, 0 replies; 9+ messages in thread
From: Chung-Lin Tang @ 2021-05-10 9:35 UTC (permalink / raw)
To: Thomas Schwinge; +Cc: gcc-patches, Catherine Moore, Tobias Burnus
On 2021/5/7 8:35 PM, Thomas Schwinge wrote:
> On 2021-05-05T23:17:25+0800, Chung-Lin Tang via Gcc-patches<gcc-patches@gcc.gnu.org> wrote:
>> This patch implements relaxing the requirements when a map with the implicit attribute encounters
>> an overlapping existing map. [...]
> Oh, oh, these data mapping interfaces/semantics ares getting more and
> more "convoluted"... %-\ (Not your fault, of course.)
>
> Haven't looked in too much detail in the patch/implementation (I'm not
> very well-versend in the exact OpenMP semantics anyway), but I suppose we
> should do similar things for OpenACC, too. I think we even currently do
> have a gimplification-level "hack" to replicate data clauses' array
> bounds for implicit data clauses on compute constructs, if the default
> "complete" mapping is going to clash with a "limited" mapping that's
> specified in an outer OpenACC 'data' directive. (That, of course,
> doesn't work for the general case of non-lexical scoping, or dynamic
> OpenACC 'enter data', etc., I suppose) I suppose your method could easily
> replace and improve that; we shall look into that later.
>
> That said, in your patch, is this current implementation (explicitly)
> meant or not meant to be active for OpenACC, too, or just OpenMP (I
> couldn't quickly tell), and/or is it (implicitly?) a no-op for OpenACC?
It appears that I have inadvertently enabled it for OpenACC as well!
But everything was tested together, so I assume it works okay for that mode as well.
The entire set of implicit-specific actions are enabled by the setting of
'OMP_CLAUSE_MAP_IMPLICIT_P (clause) = 1' inside gimplify.c:gimplify_adjust_omp_clauses_1,
so in case you want to disable it for OpenACC again, that's where you need to add the guard condition.
>> Also, another adjustment in this patch is how implicitly created clauses are added to the current
>> clause list in gimplify_adjust_omp_clauses(). Instead of simply appending the new clauses to the end,
>> this patch adds them at the position "after initial non-map clauses, but right before any existing
>> map clauses".
> Probably you haven't been testing such a configuration; I've just pushed
> "Fix up 'c-c++-common/goacc/firstprivate-mappings-1.c' for C, non-LP64"
> to devel/omp/gcc-10 branch in commit
> c51cc3b96f0b562deaffcfbcc51043aed216801a, see attached.
Thanks, I was relying on eyeballing to know where to fix testcases like this;
I did fix another similar case, but missed this one.
>
>> The reason for this is: when combined with other map clauses, for example:
>>
>> #pragma omp target map(rec.ptr[:N])
>> for (int i = 0; i < N; i++)
>> rec.ptr[i] += 1;
>>
>> There will be an implicit map created for map(rec), because of the access inside the target region.
>> The expectation is that 'rec' is implicitly mapped, and then the pointed array-section part by 'rec.ptr'
>> will be mapped, and then attachment to the 'rec.ptr' field of the mapped 'rec' (in that order).
>>
>> If the implicit 'map(rec)' is appended to the end, instead of placed before other maps, the attachment
>> operation will not find anything to attach to, and the entire region will fail.
> But that doesn't (negatively) affect user-visible semantics (OpenMP, and
> also OpenACC, if applicable), in that more/bigger objects then get mapped
> than were before? (I suppose not?)
It probably won't affect user level semantics, although we should look out if this change in convention
exposes some other bugs.
Chung-Lin
^ permalink raw reply [flat|nested] 9+ messages in thread
* [PATCH, OpenMP 5.0] Implement relaxation of implicit map vs. existing device mappings (for mainline trunk)
@ 2021-05-14 13:20 ` Chung-Lin Tang
2021-06-07 11:28 ` Thomas Schwinge
2021-06-24 15:55 ` Jakub Jelinek
0 siblings, 2 replies; 9+ messages in thread
From: Chung-Lin Tang @ 2021-05-14 13:20 UTC (permalink / raw)
To: gcc-patches, Jakub Jelinek, Catherine Moore, Tobias Burnus,
Thomas Schwinge
[-- Attachment #1: Type: text/plain, Size: 6104 bytes --]
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.
[-- Attachment #2: omp-1463.patch --]
[-- Type: text/plain, Size: 26727 bytes --]
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;
+}
^ permalink raw reply [flat|nested] 9+ messages in thread
* Re: [PATCH, OpenMP 5.0] Implement relaxation of implicit map vs. existing device mappings (for mainline trunk)
2021-05-14 13:20 ` [PATCH, OpenMP 5.0] Implement relaxation of implicit map vs. existing device mappings (for mainline trunk) Chung-Lin Tang
@ 2021-06-07 11:28 ` Thomas Schwinge
2021-06-24 15:55 ` Jakub Jelinek
1 sibling, 0 replies; 9+ messages in thread
From: Thomas Schwinge @ 2021-06-07 11:28 UTC (permalink / raw)
To: Chung-Lin Tang; +Cc: gcc-patches, Jakub Jelinek, Catherine Moore, Tobias Burnus
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
^ permalink raw reply [flat|nested] 9+ messages in thread
* Re: [PATCH, OpenMP 5.0] Implement relaxation of implicit map vs. existing device mappings (for mainline trunk)
2021-05-14 13:20 ` [PATCH, OpenMP 5.0] Implement relaxation of implicit map vs. existing device mappings (for mainline trunk) Chung-Lin Tang
2021-06-07 11:28 ` Thomas Schwinge
@ 2021-06-24 15:55 ` Jakub Jelinek
2021-11-05 16:51 ` [PATCH, v2, " Chung-Lin Tang
1 sibling, 1 reply; 9+ messages in thread
From: Jakub Jelinek @ 2021-06-24 15:55 UTC (permalink / raw)
To: Chung-Lin Tang
Cc: gcc-patches, Catherine Moore, Tobias Burnus, Thomas Schwinge
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
^ permalink raw reply [flat|nested] 9+ messages in thread
* [PATCH, v2, OpenMP 5.0] Implement relaxation of implicit map vs. existing device mappings (for mainline trunk)
2021-06-24 15:55 ` Jakub Jelinek
@ 2021-11-05 16:51 ` Chung-Lin Tang
2021-11-09 15:18 ` Jakub Jelinek
2021-11-24 11:22 ` Thomas Schwinge
0 siblings, 2 replies; 9+ messages in thread
From: Chung-Lin Tang @ 2021-11-05 16:51 UTC (permalink / raw)
To: Jakub Jelinek
Cc: gcc-patches, Catherine Moore, Tobias Burnus, Thomas Schwinge
[-- Attachment #1: Type: text/plain, Size: 10197 bytes --]
Hi Jakub,
On 2021/6/24 11:55 PM, Jakub Jelinek wrote:
> 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)?
I have changed the name to OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P, to signal that
this bit is to be passed to the runtime. Right now its intended to be used by
clauses created by the middle-end, but front-end uses like that for C++ could
be clarified later.
> 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.
Agreed, I've adjusted the patch to only to this implicit setting for OpenMP.
This reduces a lot of the originally needed scan test adjustment for existing OpenACC testcases.
>> @@ -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.
Done.
>> + 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?
As I also explained in the first submission email, due to the processing order,
if implicit classes are added last (and processed last), for example:
#pragma omp target map(tofrom: var.ptr[:N]) map(tofrom: var[implicit])
{
// access of var.ptr[]
}
The explicit var.ptr[:N] will not find anything to map, because the (implicit) map(var) has not been seen yet,
and the assumed array section attachment behavior will fail.
Only an order like: map(tofrom: var[implicit]) map(tofrom: var.ptr[:N]) will the usual assumed behavior show.
And yes, this depends on the new behavior implemented by patch [1], which I still need you to review.
e.g. for map(var.ptr[:N]), the proper behavior should *only* map the array section but NOT the base-pointer.
[1] https://gcc.gnu.org/pipermail/gcc-patches/2021-May/571195.html
> 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?).
Implicit clauses do not appear to be created if there's an explicit clause already existing.
>> +#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)
...
>> +#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.
Currently with the place of setting OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P in
gimplify_adjust_omp_clauses_1, only things like GOMP_MAP_TO/FROM/FORCE_PRESENT/etc.
be set.
I already had some trouble designating how GOMP_MAP_IMPLICIT would be encoded
within the current bits. My guess is that with more sophisticated features like
declare mappers, the whole interface probably needs to be extended somewhat.
>> @@ -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?
GOMP_MAP_IMPLICIT is (GOMP_MAP_FLAG_SPECIAL_3 | GOMP_MAP_FLAG_SPECIAL_4),
and those two bits are also partially used for some other things.
So an specific check should be needed.
> 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.
I see, I've added a short_mapkind check in get_kind.
>> + 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.
This too added.
Re-tested for trunk without regressions, is this okay now?
Thanks,
Chung-Lin
2021-11-05 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_RUNTIME_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_RUNTIME_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.
* g++.dg/goacc/firstprivate-mappings-1.C: 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.
[-- Attachment #2: implicit-20211105.patch --]
[-- Type: text/plain, Size: 17679 bytes --]
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index d8e4b139349..59e47bf2ade 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -10861,6 +10861,10 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
gcc_unreachable ();
}
OMP_CLAUSE_SET_MAP_KIND (clause, kind);
+ /* Setting of the implicit flag for the runtime is currently disabled for
+ OpenACC. */
+ if ((gimplify_omp_ctxp->region_type & ORT_ACC) == 0)
+ OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P (clause) = 1;
if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
{
@@ -11476,9 +11480,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 15e4424b0bc..3d58a6d35e6 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -13153,6 +13153,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_RUNTIME_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..74ab05bc856 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 "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 c2b8dc6c880..0a123bec58f 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_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.. map.attach:s.e .bias: 0.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_data map.attach:a .bias: 0.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_exit_data map.detach:a .bias: 0.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_data map.to:a .len: 8.." 1 "omplower" } } */
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/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 275dc7d8af7..0da85efc104 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -971,6 +971,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_RUNTIME_IMPLICIT_P (clause))
+ pp_string (pp, "[implicit]");
pp_right_paren (pp);
break;
diff --git a/gcc/tree.h b/gcc/tree.h
index 7542d97ce12..ba974471339 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1689,6 +1689,11 @@ class auto_suppress_location_wrappers
map clause. */
#define OMP_CLAUSE_MAP_IMPLICIT(NODE) \
(OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.default_def_flag)
+/* Nonzero if this map clause is to be indicated to the runtime as 'implicit',
+ due to being created through implicit data-mapping rules in the middle-end.
+ NOTE: this is different than OMP_CLAUSE_MAP_IMPLICIT. */
+#define OMP_CLAUSE_MAP_RUNTIME_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 ebd08013430..3e42d7123ae 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 196dba4f08c..dd7f573fea8 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -539,7 +539,7 @@ static inline void
gomp_map_vars_existing (struct gomp_device_descr *devicep,
struct goacc_asyncqueue *aq, splay_tree_key oldn,
splay_tree_key newn, struct target_var_desc *tgt_var,
- unsigned char kind, bool always_to_flag,
+ unsigned char kind, bool always_to_flag, bool implicit,
struct gomp_coalesce_buf *cbuf,
htab_t *refcount_set)
{
@@ -550,11 +550,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 "
@@ -564,11 +575,17 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep,
}
if (GOMP_MAP_ALWAYS_TO_P (kind) || always_to_flag)
- gomp_copy_host2dev (devicep, aq,
- (void *) (oldn->tgt->tgt_start + oldn->tgt_offset
- + newn->host_start - oldn->host_start),
- (void *) newn->host_start,
- newn->host_end - newn->host_start, false, cbuf);
+ {
+ /* Implicit + always should not happen. If this does occur, below
+ address/length adjustment is a TODO. */
+ assert (!implicit_subset);
+
+ gomp_copy_host2dev (devicep, aq,
+ (void *) (oldn->tgt->tgt_start + oldn->tgt_offset
+ + newn->host_start - oldn->host_start),
+ (void *) newn->host_start,
+ newn->host_end - newn->host_start, false, cbuf);
+ }
gomp_increment_refcount (oldn, refcount_set);
}
@@ -576,8 +593,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 (short_mapkind && 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 short_mapkind && GOMP_MAP_IMPLICIT_P (val);
}
static void
@@ -631,6 +664,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;
@@ -638,12 +672,14 @@ gomp_map_fields_existing (struct target_mem_desc *tgt,
cur_node.host_end = cur_node.host_start + sizes[i];
splay_tree_key n2 = splay_tree_lookup (mem_map, &cur_node);
kind = get_kind (short_mapkind, kinds, i);
+ implicit = get_implicit (short_mapkind, kinds, i);
if (n2
&& n2->tgt == n->tgt
&& n2->host_start - n->host_start == n2->tgt_offset - n->tgt_offset)
{
gomp_map_vars_existing (devicep, aq, n2, &cur_node, &tgt->list[i],
- kind & typemask, false, cbuf, refcount_set);
+ kind & typemask, false, implicit, cbuf,
+ refcount_set);
return;
}
if (sizes[i] == 0)
@@ -659,7 +695,8 @@ gomp_map_fields_existing (struct target_mem_desc *tgt,
== n2->tgt_offset - n->tgt_offset)
{
gomp_map_vars_existing (devicep, aq, n2, &cur_node, &tgt->list[i],
- kind & typemask, false, cbuf, refcount_set);
+ kind & typemask, false, implicit, cbuf,
+ refcount_set);
return;
}
}
@@ -671,7 +708,8 @@ gomp_map_fields_existing (struct target_mem_desc *tgt,
&& n2->host_start - n->host_start == n2->tgt_offset - n->tgt_offset)
{
gomp_map_vars_existing (devicep, aq, n2, &cur_node, &tgt->list[i],
- kind & typemask, false, cbuf, refcount_set);
+ kind & typemask, false, implicit, cbuf,
+ refcount_set);
return;
}
}
@@ -903,6 +941,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)
{
@@ -1085,8 +1124,8 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
}
}
gomp_map_vars_existing (devicep, aq, n, &cur_node, &tgt->list[i],
- kind & typemask, always_to_cnt > 0, NULL,
- refcount_set);
+ kind & typemask, always_to_cnt > 0, implicit,
+ NULL, refcount_set);
i += always_to_cnt;
}
else
@@ -1256,6 +1295,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)
@@ -1415,7 +1455,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
splay_tree_key n = splay_tree_lookup (mem_map, k);
if (n && n->refcount != REFCOUNT_LINK)
gomp_map_vars_existing (devicep, aq, n, k, &tgt->list[i],
- kind & typemask, false, cbufp,
+ kind & typemask, false, implicit, cbufp,
refcount_set);
else
{
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;
+}
^ permalink raw reply [flat|nested] 9+ messages in thread
* Re: [PATCH, v2, OpenMP 5.0] Implement relaxation of implicit map vs. existing device mappings (for mainline trunk)
2021-11-05 16:51 ` [PATCH, v2, " Chung-Lin Tang
@ 2021-11-09 15:18 ` Jakub Jelinek
2021-11-24 11:22 ` Thomas Schwinge
1 sibling, 0 replies; 9+ messages in thread
From: Jakub Jelinek @ 2021-11-09 15:18 UTC (permalink / raw)
To: Chung-Lin Tang
Cc: gcc-patches, Catherine Moore, Tobias Burnus, Thomas Schwinge
On Sat, Nov 06, 2021 at 12:51:59AM +0800, Chung-Lin Tang wrote:
> 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 (short_mapkind && GOMP_MAP_IMPLICIT_P (val))
> + val &= ~GOMP_MAP_IMPLICIT;
> + return val;
> +}
It doesn't make sense to test it twice. I'd do:
if (!short_mapkind)
return ((unsigned char *) kinds)[idx];
int val = ((unsigned short *) 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 short_mapkind && GOMP_MAP_IMPLICIT_P (val);
> }
and here even simpler, no need to read kinds at all:
if (!short_mapkind)
return false;
int val = ((unsigned short *) kinds)[idx];
return GOMP_MAP_IMPLICIT_P (val);
?
Otherwise LGTM.
Jakub
^ permalink raw reply [flat|nested] 9+ messages in thread
* Re: [PATCH, v2, OpenMP 5.0] Implement relaxation of implicit map vs. existing device mappings (for mainline trunk)
2021-11-05 16:51 ` [PATCH, v2, " Chung-Lin Tang
2021-11-09 15:18 ` Jakub Jelinek
@ 2021-11-24 11:22 ` Thomas Schwinge
1 sibling, 0 replies; 9+ messages in thread
From: Thomas Schwinge @ 2021-11-24 11:22 UTC (permalink / raw)
To: Chung-Lin Tang, gcc-patches; +Cc: Catherine Moore, Tobias Burnus, Jakub Jelinek
[-- Attachment #1: Type: text/plain, Size: 5318 bytes --]
Hi!
On 2021-11-06T00:51:59+0800, Chung-Lin Tang <cltang@codesourcery.com> wrote:
> On 2021/6/24 11:55 PM, Jakub Jelinek wrote:
>> 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)
>>> {
>> Also as Thomas mentioned, it should be restricted to non-OpenACC,
> Agreed, I've adjusted the patch to only to this implicit setting for OpenMP.
> This reduces a lot of the originally needed scan test adjustment for existing OpenACC testcases.
..., but not all, because this piece is still effective:
>>> @@ -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);
..., which effects changes such as:
> --- a/gcc/testsuite/c-c++-common/goacc/combined-reduction.c
> +++ b/gcc/testsuite/c-c++-common/goacc/combined-reduction.c
> -/* { 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" } } */
> --- a/gcc/testsuite/c-c++-common/goacc/mdc-1.c
> +++ b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
> -/* { 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.. map.attach:s.e .bias: 0.." 1 "omplower" } } */
> --- a/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C
> +++ b/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C
> - { 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 } }
..., and you've changed:
> --- 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;
> }
..., however the clauses reordering alone isn't going to fix PR90859
"[OMP] Mappings for VLA different depending on 'target { c && { ! lp64 } }'",
so it's not correct to just remove that testing/documentation here -- this
change gave rise to PR103244
"c-c++-common/goacc/firstprivate-mappings-1.c fails in certain
configurations since g:b7e20480630e3eeb9eed8b3941da3b3f0c22c969". To
resolve that, and until we properly and deliberately look into also for
OpenACC enabling your "Implement relaxation of implicit map vs. existing
device mappings" (we certainly should!), I've now pushed to master branch
commit fdd34569e7a9fc2b6c638a7ef62b965ed7e832ce "Restore previous OpenACC
implicit data clauses ordering [PR103244]", see attached.
Grüße
Thomas
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-Restore-previous-OpenACC-implicit-data-clauses-order.patch --]
[-- Type: text/x-diff, Size: 7076 bytes --]
From fdd34569e7a9fc2b6c638a7ef62b965ed7e832ce Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Mon, 22 Nov 2021 17:29:09 +0100
Subject: [PATCH] Restore previous OpenACC implicit data clauses ordering
[PR103244]
Follow-up for recent commit b7e20480630e3eeb9eed8b3941da3b3f0c22c969
"openmp: Relax handling of implicit map vs. existing device mappings".
As discussed, we likely also for OpenACC ought to use
'OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P' and do the appropriate implicit clauses
ordering -- but that's for a separate step.
gcc/
PR middle-end/103244
* gimplify.c (gimplify_adjust_omp_clauses): Restore previous
OpenACC behavior.
gcc/testsuite/
PR middle-end/103244
* c-c++-common/goacc/combined-reduction.c: Revert/expect previous
OpenACC behavior.
* c-c++-common/goacc/firstprivate-mappings-1.c: Likewise.
* c-c++-common/goacc/mdc-1.c: Likewise.
* g++.dg/goacc/firstprivate-mappings-1.C: Likewise.
---
gcc/gimplify.c | 22 ++++++++++++-------
.../c-c++-common/goacc/combined-reduction.c | 2 +-
.../goacc/firstprivate-mappings-1.c | 7 +++++-
gcc/testsuite/c-c++-common/goacc/mdc-1.c | 2 +-
.../g++.dg/goacc/firstprivate-mappings-1.C | 2 +-
5 files changed, 23 insertions(+), 12 deletions(-)
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 816cdaf8a18..8624f8221fd 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -11501,15 +11501,21 @@ 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. Implicit clauses are added at the start
- of the clause list, but after any non-map clauses. */
+ /* Add in any implicit data sharing. */
struct gimplify_adjust_omp_clauses_data data;
- 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;
+ if ((gimplify_omp_ctxp->region_type & ORT_ACC) == 0)
+ {
+ /* OpenMP. Implicit clauses are added at the start of the clause list,
+ but after any non-map clauses. */
+ 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;
+ }
+ else
+ /* OpenACC. */
+ data.list_p = list_p;
data.pre_p = pre_p;
splay_tree_foreach (ctx->variables, gimplify_adjust_omp_clauses_1, &data);
diff --git a/gcc/testsuite/c-c++-common/goacc/combined-reduction.c b/gcc/testsuite/c-c++-common/goacc/combined-reduction.c
index 74ab05bc856..ecf23f59d66 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. firstprivate.n. map.tofrom:v1" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "omp target oacc_parallel reduction.+:v1. 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 "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 5134ef6ed6c..7987beaed9a 100644
--- a/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/firstprivate-mappings-1.c
@@ -419,7 +419,12 @@ vla (int array_li)
copyout (array_so)
/* The gimplifier has created an implicit 'firstprivate' clause for the array
length.
- { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel firstprivate\(array_li.[0-9]+\) map\(from:array_so \[len: 4\]\) \[} omplower } } */
+ { 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 } } } } } */
{
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 0a123bec58f..c2b8dc6c880 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_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.tofrom:s .len: 32.. map.attach:s.e .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_enter_data map.attach:a .bias: 0.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_exit_data map.detach:a .bias: 0.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_data map.to:a .len: 8.." 1 "omplower" } } */
diff --git a/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C b/gcc/testsuite/g++.dg/goacc/firstprivate-mappings-1.C
index 99a3bd472f7..1b1badb1a90 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 firstprivate\([^)]+\) map\(from:array_so \[len: 4\]\)} omplower } }
+ { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_parallel map\(from:array_so \[len: 4\]\) firstprivate\(} omplower } }
(C++ computes an intermediate value, so can't scan for 'firstprivate(array_li)'.) */
{
array_so = sizeof array;
--
2.33.0
^ permalink raw reply [flat|nested] 9+ messages in thread
end of thread, other threads:[~2021-11-24 11:22 UTC | newest]
Thread overview: 9+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2021-05-05 15:17 [PATCH, OG10, OpenMP 5.0, committed] Implement relaxation of implicit map vs. existing device mappings Chung-Lin Tang
2021-05-07 12:35 ` Thomas Schwinge
2021-05-10 9:35 ` Chung-Lin Tang
2021-05-14 13:20 ` [PATCH, OpenMP 5.0] Implement relaxation of implicit map vs. existing device mappings (for mainline trunk) Chung-Lin Tang
2021-06-07 11:28 ` Thomas Schwinge
2021-06-24 15:55 ` Jakub Jelinek
2021-11-05 16:51 ` [PATCH, v2, " Chung-Lin Tang
2021-11-09 15:18 ` Jakub Jelinek
2021-11-24 11:22 ` Thomas Schwinge
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).