diff mbox series

[14/14] OpenACC: Improve implicit mapping for non-lexically nested offload regions

Message ID a02f67664b5ffd7c70dc1cd251673e40fbc2b569.1687201316.git.julian@codesourcery.com
State New
Headers show
Series OpenMP/OpenACC: map clause and OMP gimplify rework | expand

Commit Message

Julian Brown June 19, 2023, 9:17 p.m. UTC 
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.
---
 gcc/gimplify.cc                               |  5 +---
 .../c-c++-common/goacc/combined-reduction.c   |  2 +-
 .../c-c++-common/goacc/reduction-1.c          |  4 ++--
 .../c-c++-common/goacc/reduction-10.c         |  9 +++----
 .../c-c++-common/goacc/reduction-2.c          |  4 ++--
 .../c-c++-common/goacc/reduction-3.c          |  4 ++--
 .../c-c++-common/goacc/reduction-4.c          |  4 ++--
 .../gfortran.dg/goacc/loop-tree-1.f90         |  2 +-
 include/gomp-constants.h                      |  3 ++-
 .../implicit-mapping-1.c                      | 24 +++++++++++++++++++
 10 files changed, 42 insertions(+), 19 deletions(-)
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/implicit-mapping-1.c
diff mbox series

Patch

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/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/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/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;
+}