As flagged by Thomas, the GOMP_MAP_STRUCT mapping is not itself necessary
for OpenACC "exit data" directives, and is skipped over (now) in libgomp.
We might as well not emit it to start with, in line with the equivalent
OpenMP directive. We can keep the "no-op" handling in libgomp for the
reason of backward compatibility.
I've added a new scan test for the new behaviour. OK?
Julian
gcc/
* gimplify.c (gimplify_adjust_omp_clauses): Remove GOMP_MAP_STRUCT
mapping from OpenACC exit data directives.
---
gcc/gimplify.c | 3 +-
.../goacc/struct-enter-exit-data-1.c | 35 +++++++++++++++++++
2 files changed, 37 insertions(+), 1 deletion(-)
create mode 100644 gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index cb08b26dc65..e14932fafaf 100644
--- 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
index 00000000000..4be5674b23e
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c
@@ -0,0 +1,35 @@
+/* Check that extraneous GOMP_MAP_STRUCTs are removed from OpenACC exit data
+ directives. */
+
+/* { dg-additional-options "-fdump-tree-omplower" } */
+
+#include <stdlib.h>
+
+struct str {
+ int a;
+ int *b;
+ int *c;
+ int d;
+};
+
+#define N 1024
+
+int
+main (int argc, char *argv[])
+{
+ struct str s;
+
+ s.b = (int *) malloc (sizeof (int) * N);
+ s.c = (int *) malloc (sizeof (int) * N);
+
+ #pragma acc enter data copyin(s.a, s.b[0:N], s.c[0:N], s.d)
+ /* { 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.d \[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\]\)} omplower } }
*/
+
+ #pragma acc exit data copyout(s.a, s.b[0:N], s.c[0:N], s.d)
+ /* { 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.d \[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\]\)} omplower }
} */
+
+ free (s.b);
+ free (s.c);
+
+ return 0;
+}
--
2.23.0
