From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1410) id 9BE19385841B; Mon, 19 Jun 2023 22:18:03 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 9BE19385841B DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gcc.gnu.org; s=default; t=1687213083; bh=Zn/875glaS4WGjMZrHcT3fxndgXC4xBtXCj3s0Ptsxk=; h=From:To:Subject:Date:From; b=oWReyPctc9BeLHuT6LpCEURpcAVGaCPR6k1+oV0hGDs933cZH1bf1kHXE1aORBsK6 lWowWIBBQEy5zQ59X0Idh/yZBMtrl1JLjCFS7iVapyKtrxn7fBOnTfBWfsyHnWNZvZ 6f0auh7CwEBK9rZ+1lYLEVaFG/uP20qCI0m2A4FI= Content-Type: text/plain; charset="us-ascii" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit From: Julian Brown To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/omp/gcc-13] OpenACC: Improve implicit mapping for non-lexically nested offload regions X-Act-Checkin: gcc X-Git-Author: Julian Brown X-Git-Refname: refs/heads/devel/omp/gcc-13 X-Git-Oldrev: 7da2c5758d263348105fc7fa416c7f1c2ca781a4 X-Git-Newrev: 29bc958efc8798b284b816174cc821f188e1e937 Message-Id: <20230619221803.9BE19385841B@sourceware.org> Date: Mon, 19 Jun 2023 22:18:03 +0000 (GMT) List-Id: https://gcc.gnu.org/g:29bc958efc8798b284b816174cc821f188e1e937 commit 29bc958efc8798b284b816174cc821f188e1e937 Author: Julian Brown Date: Sat Jun 10 15:17:19 2023 +0000 OpenACC: Improve implicit mapping for non-lexically nested offload regions This patch enables use of the OMP_CLAUSE_RUNTIME_IMPLICIT_P flag for OpenACC. This allows code like this to work correctly: int arr[100]; [...] #pragma acc enter data copyin(arr[20:10]) /* No explicit mapping of 'arr' here. */ #pragma acc parallel { /* use of arr[20:10]... */ } #pragma acc exit data copyout(arr[20:10]) Otherwise, the implicit "copy" ("present_or_copy") on the parallel corresponds to the whole array, and that fails at runtime when the subarray is mapped. The numbering of the GOMP_MAP_IMPLICIT bit clashes with the OpenACC "non-contiguous" dynamic array support, so the GOMP_MAP_NONCONTIG_ARRAY_P macro has been adjusted to account for that. This behaviour relates to upstream OpenACC issue 490 (not yet resolved). 2023-06-16 Julian Brown gcc/ * gimplify.cc (gimplify_adjust_omp_clauses_1): Set OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P for OpenACC also. gcc/testsuite/ * c-c++-common/goacc/combined-reduction.c: Adjust scan output. * 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-10.c: Likewise. * gfortran.dg/goacc/loop-tree-1.f90: Likewise. include/ * gomp-constants.h (GOMP_MAP_NONCONTIG_ARRAY_P): Tweak condition. libgomp/ * testsuite/libgomp.oacc-c-c++-common/implicit-mapping-1.c: New test. Diff: --- gcc/ChangeLog.omp | 5 +++++ gcc/gimplify.cc | 5 +---- gcc/testsuite/ChangeLog.omp | 10 +++++++++ .../c-c++-common/goacc/combined-reduction.c | 2 +- gcc/testsuite/c-c++-common/goacc/reduction-1.c | 4 ++-- gcc/testsuite/c-c++-common/goacc/reduction-10.c | 9 ++++---- 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/gfortran.dg/goacc/loop-tree-1.f90 | 2 +- include/ChangeLog.omp | 4 ++++ include/gomp-constants.h | 3 ++- libgomp/ChangeLog.omp | 4 ++++ .../libgomp.oacc-c-c++-common/implicit-mapping-1.c | 24 ++++++++++++++++++++++ 14 files changed, 65 insertions(+), 19 deletions(-) diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index ca6ec455137..cbc29c12ca2 100644 --- a/gcc/ChangeLog.omp +++ b/gcc/ChangeLog.omp @@ -1,3 +1,8 @@ +2023-06-19 Julian Brown + + * gimplify.cc (gimplify_adjust_omp_clauses_1): Set + OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P for OpenACC also. + 2023-06-19 Julian Brown * gimplify.cc (gimplify_adjust_omp_clauses_1): Handle "oacc declare diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 0706f130ebb..1e90d2ed031 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -13413,10 +13413,7 @@ 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; + OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P (clause) = 1; if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST) { diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp index 709d85f7072..ba530d3f6c4 100644 --- a/gcc/testsuite/ChangeLog.omp +++ b/gcc/testsuite/ChangeLog.omp @@ -1,3 +1,13 @@ +2023-06-19 Julian Brown + + * c-c++-common/goacc/combined-reduction.c: Adjust scan output. + * 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-10.c: Likewise. + * gfortran.dg/goacc/loop-tree-1.f90: Likewise. + 2023-06-19 Julian Brown * gfortran.dg/goacc/assumed-size.f90: Don't expect error. diff --git a/gcc/testsuite/c-c++-common/goacc/combined-reduction.c b/gcc/testsuite/c-c++-common/goacc/combined-reduction.c index ecf23f59d66..40b93acc9ea 100644 --- a/gcc/testsuite/c-c++-common/goacc/combined-reduction.c +++ b/gcc/testsuite/c-c++-common/goacc/combined-reduction.c @@ -25,5 +25,5 @@ main () /* { 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 "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/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-10.c b/gcc/testsuite/c-c++-common/goacc/reduction-10.c index 579aa561479..36c330e9267 100644 --- a/gcc/testsuite/c-c++-common/goacc/reduction-10.c +++ b/gcc/testsuite/c-c++-common/goacc/reduction-10.c @@ -87,7 +87,8 @@ 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\]\)} 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:result .len: 4.. map.tofrom:array .len: 4000..implicit.." 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/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/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/include/ChangeLog.omp b/include/ChangeLog.omp index bfcd25d8b7f..2b0a2211e92 100644 --- a/include/ChangeLog.omp +++ b/include/ChangeLog.omp @@ -1,3 +1,7 @@ +2023-06-19 Julian Brown + + * gomp-constants.h (GOMP_MAP_NONCONTIG_ARRAY_P): Tweak condition. + 2023-06-12 Tobias Burnus Backported from mainline: diff --git a/include/gomp-constants.h b/include/gomp-constants.h index b8281b81800..0f8f0f31f4e 100644 --- a/include/gomp-constants.h +++ b/include/gomp-constants.h @@ -276,7 +276,8 @@ enum gomp_map_kind || (X) == GOMP_MAP_FORCE_PRESENT) #define GOMP_MAP_NONCONTIG_ARRAY_P(X) \ - ((X) & GOMP_MAP_NONCONTIG_ARRAY) + (((X) & GOMP_MAP_NONCONTIG_ARRAY) != 0 \ + && ((X) & GOMP_MAP_FLAG_SPECIAL_4) == 0) /* 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 cb84dee3601..39697342010 100644 --- a/libgomp/ChangeLog.omp +++ b/libgomp/ChangeLog.omp @@ -1,3 +1,7 @@ +2023-06-19 Julian Brown + + * testsuite/libgomp.oacc-c-c++-common/implicit-mapping-1.c: New test. + 2023-06-19 Julian Brown * testsuite/libgomp.oacc-fortran/nonlexical-assumed-size-1.f90: New diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/implicit-mapping-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/implicit-mapping-1.c new file mode 100644 index 00000000000..4825e875998 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/implicit-mapping-1.c @@ -0,0 +1,24 @@ +/* { dg-do run } */ + +#include +#include + +int main(void) +{ + int arr[100]; + + memset (arr, 0, sizeof (int) * 100); + +#pragma acc enter data copyin(arr[30:10]) + +#pragma acc serial + { + arr[33] = 66; + } + +#pragma acc exit data copyout(arr[30:10]) + + assert (arr[33] == 66); + + return 0; +}