diff mbox series

[2/3] Strip GOMP_MAP_STRUCT from OpenACC exit data mappings

Message ID cf6237ce3f7cad564e2581e450c7d0346f53365f.1591276990.git.julian@codesourcery.com
State New
Headers show
Series OpenACC "exit data" copyout, and Fortran derived-type members | expand

Commit Message

Julian Brown June 4, 2020, 1:40 p.m. UTC 
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 mbox series

Patch

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