[OpenACC 'exit data'] Strip 'GOMP_MAP_STRUCT' mappings

These are not itself necessary for OpenACC 'exit data' directives, and are
skipped over (now) in libgomp.  We might as well not emit them to start with,
in line with the equivalent OpenMP directive.  We keep the no-op handling in
libgomp for the reason of backward compatibility.

	gcc/
	* gimplify.c (gimplify_adjust_omp_clauses): Remove
	'GOMP_MAP_STRUCT' mapping from OpenACC 'exit data' directives.
	gcc/testsuite/
	* c-c++-common/goacc/struct-enter-exit-data-1.c: New file.
	libgomp/
	* oacc-mem.c (goacc_exit_data_internal) <GOMP_MAP_STRUCT>: Explain
	special handling.

Co-Authored-By: Julian Brown <julian@codesourcery.com>

diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index cb08b26dc65e7be0706f4eb1dc4a0de096c5c5c3..e14932fafaf8e7a39e1069cf15b0ab4b66783a54 100644 (file)
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -10396,7 +10396,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
                }
            }
          else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT
-                  && code == OMP_TARGET_EXIT_DATA)
+                  && (code == OMP_TARGET_EXIT_DATA
+                      || code == OACC_EXIT_DATA))
            remove = true;
          else if (DECL_SIZE (decl)
                   && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST

diff --git a/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c b/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c
new file mode 100644 (file)
index 0000000..df405e4
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c
@@ -0,0 +1,27 @@
+/* Check 'GOMP_MAP_STRUCT' mapping, and in particular that it gets removed from
+   OpenACC 'exit data' directives.  */
+
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+struct str {
+  int a;
+  int *b;
+  int *c;
+  int d;
+  int *e;
+  int f;
+};
+
+#define N 1024
+
+void
+test (int *b, int *c, int *e)
+{
+  struct str s = { .a = 0, .b = b, .c = c, .d = 0, .e = e, .f = 0 };
+
+#pragma acc enter data copyin(s.a, s.b[0:N], s.c[0:N] /* , s.d */ /* , s.e[0:N] */, s.f)
+  /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(struct:s \[len: 4\]\) map\(to:s.a \[len: [0-9]+\]\) map\(alloc:s.b \[len: [0-9]+\]\) map\(alloc:s.c \[len: [0-9]+\]\) map\(to:s.f \[len: [0-9]+\]\) map\(to:\*[_0-9]+ \[len: [0-9]+\]\) map\(attach:s.b \[bias: 0\]\) map\(to:\*[_0-9]+ \[len: [0-9]+\]\) map\(attach:s.c \[bias: 0\]\)$} gimple } } */
+
+#pragma acc exit data copyout(s.a, s.b[0:N], s.c[0:N] /* , s.d */ /* , s.e[0:N] */, s.f)
+  /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(from:s.a \[len: [0-9]+\]\) map\(release:s.b \[len: [0-9]+\]\) map\(release:s.c \[len: [0-9]+\]\) map\(from:s.f \[len: [0-9]+\]\) map\(from:\*[_0-9]+ \[len: [0-9]+\]\) map\(detach:s.b \[bias: 0\]\) map\(from:\*[_0-9]+ \[len: [0-9]+\]\) map\(detach:s.c \[bias: 0\]\)$} gimple } } */
+}

diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 1e3685a073da935d677369b7346966159900e936..936ae649dd934be6403705a3a22b337d74f06440 100644 (file)
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -1181,8 +1181,9 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
 
        case GOMP_MAP_STRUCT:
          /* Skip the 'GOMP_MAP_STRUCT' itself, and use the regular processing
-            for all its entries.  TODO: don't generate these no-op
-            'GOMP_MAP_STRUCT's.  */
+            for all its entries.  This special handling exists for GCC 10.1
+            compatibility; afterwards, we're not generating these no-op
+            'GOMP_MAP_STRUCT's anymore.  */
          break;
 
        default: