public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-13] OpenACC: Improve implicit mapping for non-lexically nested offload regions
@ 2023-06-19 22:18 Julian Brown
0 siblings, 0 replies; only message in thread
From: Julian Brown @ 2023-06-19 22:18 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:29bc958efc8798b284b816174cc821f188e1e937
commit 29bc958efc8798b284b816174cc821f188e1e937
Author: Julian Brown <julian@codesourcery.com>
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 <julian@codesourcery.com>
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 <julian@codesourcery.com>
+
+ * gimplify.cc (gimplify_adjust_omp_clauses_1): Set
+ OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P for OpenACC also.
+
2023-06-19 Julian Brown <julian@codesourcery.com>
* 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 <julian@codesourcery.com>
+
+ * 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 <julian@codesourcery.com>
* 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 <julian@codesourcery.com>
+
+ * gomp-constants.h (GOMP_MAP_NONCONTIG_ARRAY_P): Tweak condition.
+
2023-06-12 Tobias Burnus <tobias@codesourcery.com>
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 <julian@codesourcery.com>
+
+ * testsuite/libgomp.oacc-c-c++-common/implicit-mapping-1.c: New test.
+
2023-06-19 Julian Brown <julian@codesourcery.com>
* 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 <string.h>
+#include <assert.h>
+
+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;
+}
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2023-06-19 22:18 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2023-06-19 22:18 [gcc/devel/omp/gcc-13] OpenACC: Improve implicit mapping for non-lexically nested offload regions Julian Brown
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).