This patch removes unnecessary special-case code in oacc-mem.c
which prevented copy-back of structure members on OpenACC "exit data"
directives. I've added a couple of new testcases to verify the behaviour
(which fail without the patch).
(On editing this mail, I notice I omitted to add a comment about
GOMP_MAP_STRUCT being a no-op, and/or no longer emitted for OpenACC
exit data with the patch later in this series. I'll add such a comment
before commit.)
OK?
Thanks,
Julian
libgomp/
* oacc-mem.c (goacc_exit_data_internal): Remove special-case
(no-copyback) handling for GOMP_MAP_STRUCT.
* testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c: New test.
* testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c: New test.
---
libgomp/oacc-mem.c | 32 --------------
.../struct-copyout-1.c | 38 ++++++++++++++++
.../struct-copyout-2.c | 44 +++++++++++++++++++
3 files changed, 82 insertions(+), 32 deletions(-)
create mode 100644
libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c
create mode 100644
libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 2d4bba78efd..232683a85f0 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -1136,38 +1136,6 @@ goacc_exit_data_internal (struct gomp_device_descr
*acc_dev, size_t mapnum,
break;
case GOMP_MAP_STRUCT:
- {
- int elems = sizes[i];
- for (int j = 1; j <= elems; j++)
- {
- struct splay_tree_key_s k;
- k.host_start = (uintptr_t) hostaddrs[i + j];
- k.host_end = k.host_start + sizes[i + j];
- splay_tree_key str;
- str = splay_tree_lookup (&acc_dev->mem_map, &k);
- if (str)
- {
- if (finalize)
- {
- if (str->refcount != REFCOUNT_INFINITY)
- str->refcount -= str->virtual_refcount;
- str->virtual_refcount = 0;
- }
- if (str->virtual_refcount > 0)
- {
- if (str->refcount != REFCOUNT_INFINITY)
- str->refcount--;
- str->virtual_refcount--;
- }
- else if (str->refcount > 0
- && str->refcount != REFCOUNT_INFINITY)
- str->refcount--;
- if (str->refcount == 0)
- gomp_remove_var_async (acc_dev, str, aq);
- }
- }
- i += elems;
- }
break;
default:
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c
b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c
new file mode 100644
index 00000000000..b86f1c921a9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c
@@ -0,0 +1,38 @@
+#include <assert.h>
+
+struct str1 {
+ int a;
+ int b;
+};
+
+struct str2 {
+ int c;
+ int d;
+ struct str1 s;
+};
+
+int
+main (int argc, char *argv[])
+{
+ struct str2 t;
+
+ t.c = 1;
+ t.d = 2;
+ t.s.a = 3;
+ t.s.b = 4;
+
+ #pragma acc enter data copyin(t.s)
+
+ #pragma acc serial present(t.s) /* { dg-warning "using vector_length
\\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */
+ {
+ t.s.a = 5;
+ t.s.b = 6;
+ }
+
+ #pragma acc exit data copyout(t.s)
+
+ assert (t.s.a == 5);
+ assert (t.s.b == 6);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c
b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c
new file mode 100644
index 00000000000..4dd8a3a7e17
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c
@@ -0,0 +1,44 @@
+#include <assert.h>
+#include <stdlib.h>
+
+struct str1 {
+ int a;
+ int b;
+ int *c;
+};
+
+#define N 1024
+
+int
+main (int argc, char *argv[])
+{
+ struct str1 s;
+
+ s.a = 1;
+ s.b = 2;
+ s.c = (int *) malloc (sizeof (int) * N);
+
+ for (int i = 0; i < N; i++)
+ s.c[i] = i + 10;
+
+ #pragma acc enter data copyin(s.a, s.b, s.c[0:N])
+
+ #pragma acc serial present(s.a, s.b, s.c[0:N]) /* { dg-warning "using
vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected }
} */
+ {
+ s.a = 3;
+ s.b = 4;
+ for (int i = 0; i < N; i++)
+ s.c[i] = i + 20;
+ }
+
+ #pragma acc exit data copyout(s.a, s.b, s.c[0:N])
+
+ assert (s.a == 3);
+ assert (s.b == 4);
+ for (int i = 0; i < N; i++)
+ assert (s.c[i] == i + 20);
+
+ free (s.c);
+
+ return 0;
+}
--
2.23.0
