diff mbox series

[1/3] OpenACC "exit data" copyout for struct members

Message ID 6538c388d22e016bd01737e8c2f80437f432c70f.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 
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 mbox series

Patch

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