From 27e985479e931771472663cad34f8b99c6f62627 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Wed, 20 May 2020 10:56:55 +0200
Subject: [PATCH] [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>
(cherry picked from commit 1afc4672561a41dfbf4e3f2c1f35f7a5b7a20339)
---
gcc/gimplify.c | 3 ++-
.../goacc/struct-enter-exit-data-1.c | 27 +++++++++++++++++++
libgomp/oacc-mem.c | 5 ++--
3 files changed, 32 insertions(+), 3 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c
@@ -10408,7 +10408,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
new file mode 100644
@@ -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 } } */
+}
@@ -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:
--
2.26.2