From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1729) id A00503AA8033; Thu, 13 May 2021 16:21:28 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org A00503AA8033 Content-Type: text/plain; charset="us-ascii" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit From: Kwok Yeung To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/omp/gcc-11] OpenMP 5.0: Implement relaxation of implicit map vs. existing device mappings X-Act-Checkin: gcc X-Git-Author: Chung-Lin Tang X-Git-Refname: refs/heads/devel/omp/gcc-11 X-Git-Oldrev: f290996b035c5739dde57e5e4a7f39b3c739b5c0 X-Git-Newrev: 9a4a899134779cef4553f0c616e9577b379783ac Message-Id: <20210513162128.A00503AA8033@sourceware.org> Date: Thu, 13 May 2021 16:21:28 +0000 (GMT) X-BeenThere: gcc-cvs@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-cvs mailing list List-Unsubscribe: , List-Archive: List-Help: List-Subscribe: , X-List-Received-Date: Thu, 13 May 2021 16:21:28 -0000 https://gcc.gnu.org/g:9a4a899134779cef4553f0c616e9577b379783ac commit 9a4a899134779cef4553f0c616e9577b379783ac Author: Chung-Lin Tang Date: Wed May 5 08:11:19 2021 -0700 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 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. Diff: --- gcc/ChangeLog.omp | 13 ++++ gcc/gimplify.c | 11 ++- gcc/omp-low.c | 13 ++++ gcc/testsuite/ChangeLog.omp | 20 ++++++ .../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/ChangeLog.omp | 9 +++ include/gomp-constants.h | 17 ++++- libgomp/ChangeLog.omp | 12 ++++ libgomp/target.c | 78 ++++++++++++++++------ .../libgomp.c-c++-common/target-implicit-map-1.c | 31 +++++++++ 28 files changed, 265 insertions(+), 57 deletions(-) diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index dc5a0321e85..add9b1754fa 100644 --- a/gcc/ChangeLog.omp +++ b/gcc/ChangeLog.omp @@ -1,3 +1,16 @@ +2021-05-05 Chung-Lin Tang + + * 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. + 2021-04-30 Kwok Cheung Yeung Backport from mainline diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 0dd8e7d71b4..e49e543fc35 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -10670,6 +10670,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) { @@ -11286,9 +11287,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 0251dbe7186..a8675800c56 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -13021,6 +13021,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/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp index 9a3e21a579b..89fc9619722 100644 --- a/gcc/testsuite/ChangeLog.omp +++ b/gcc/testsuite/ChangeLog.omp @@ -1,3 +1,23 @@ +2021-05-05 Chung-Lin Tang + + * 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. + 2021-04-30 Kwok Cheung Yeung * c-c++-common/goacc/note-parallelism-1-kernels-straight-line.c: Add 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..f43e4b46cb6 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 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/goacc/reduction-8.c b/gcc/testsuite/c-c++-common/goacc/reduction-8.c index 8a0283f4ac3..8494e594665 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 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/g++.dg/gomp/target-lambda-1.C b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C index 7dceef80f47..e5a24d7abc4 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 08568f9284c..2755b4b58bd 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 3b2d5811350..3703762f45a 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 e43d376fc28..e9f169f9517 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 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 0c47045df9c..fef512612bd 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 3357a20263e..38459cfadf3 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 18f4c622180..56d0fa6bd6e 100644 --- a/gcc/tree-pretty-print.c +++ b/gcc/tree-pretty-print.c @@ -977,6 +977,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 edf4fd37979..72cce9778f2 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1649,6 +1649,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/ChangeLog.omp b/include/ChangeLog.omp index a460e639d22..c4b3e729533 100644 --- a/include/ChangeLog.omp +++ b/include/ChangeLog.omp @@ -1,3 +1,12 @@ +2021-05-05 Chung-Lin Tang + + * 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. + 2018-10-04 Cesar Philippidis Julian Brown diff --git a/include/gomp-constants.h b/include/gomp-constants.h index d9b80ac9732..a995ad9379f 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/ChangeLog.omp b/libgomp/ChangeLog.omp index b605e4b8fe7..6a4a656624a 100644 --- a/libgomp/ChangeLog.omp +++ b/libgomp/ChangeLog.omp @@ -1,3 +1,15 @@ +2021-05-05 Chung-Lin Tang + + * 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. + 2021-04-30 Kwok Cheung Yeung Backport from mainline diff --git a/libgomp/target.c b/libgomp/target.c index 2633ab1e43f..7a8fa87bd85 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 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; +}