OpenACC 2.6 deep copy: C and C++ execution tests

libgomp/
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-1.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-2.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-4.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-6.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-9.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-11.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-14.c: New test.
	* testsuite/libgomp.oacc-c++/deep-copy-12.C: New test.
	* testsuite/libgomp.oacc-c++/deep-copy-13.C: New test.

From-SVN: r279629

libgomp/ChangeLog

@@ -1,3 +1,18 @@
+2019-12-19  Julian Brown  <julian@codesourcery.com>
+
+	* testsuite/libgomp.oacc-c-c++-common/deep-copy-1.c: New test.
+	* testsuite/libgomp.oacc-c-c++-common/deep-copy-2.c: New test.
+	* testsuite/libgomp.oacc-c-c++-common/deep-copy-4.c: New test.
+	* testsuite/libgomp.oacc-c-c++-common/deep-copy-6.c: New test.
+	* testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c: New test.
+	* testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c: New test.
+	* testsuite/libgomp.oacc-c-c++-common/deep-copy-9.c: New test.
+	* testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c: New test.
+	* testsuite/libgomp.oacc-c-c++-common/deep-copy-11.c: New test.
+	* testsuite/libgomp.oacc-c-c++-common/deep-copy-14.c: New test.
+	* testsuite/libgomp.oacc-c++/deep-copy-12.C: New test.
+	* testsuite/libgomp.oacc-c++/deep-copy-13.C: New test.
+
 2019-12-19  Julian Brown  <julian@codesourcery.com>
 
 	* libgomp.h (struct target_var_desc): Add do_detach flag.

libgomp/testsuite/libgomp.oacc-c++/deep-copy-12.C

@@ -0,0 +1,72 @@
+#include <stdlib.h>
+
+/* Test attach/detach with dereferences of reference to pointer to struct.  */
+
+typedef struct {
+  int *a;
+  int *b;
+  int *c;
+} mystruct;
+
+int main(int argc, char* argv[])
+{
+  const int N = 1024;
+  mystruct *m = (mystruct *) malloc (sizeof (*m));
+  mystruct *&mref = m;
+  int i;
+
+  mref->a = (int *) malloc (N * sizeof (int));
+  m->b = (int *) malloc (N * sizeof (int));
+  m->c = (int *) malloc (N * sizeof (int));
+
+  for (i = 0; i < N; i++)
+    {
+      mref->a[i] = 0;
+      m->b[i] = 0;
+      m->c[i] = 0;
+    }
+
+#pragma acc enter data copyin(m[0:1])
+
+  for (int i = 0; i < 99; i++)
+    {
+      int j;
+#pragma acc parallel loop copy(mref->a[0:N])
+      for (j = 0; j < N; j++)
+	mref->a[j]++;
+#pragma acc parallel loop copy(mref->b[0:N], m->c[5:N-10])
+      for (j = 0; j < N; j++)
+	{
+	  mref->b[j]++;
+	  if (j > 5 && j < N - 5)
+	    m->c[j]++;
+	}
+    }
+
+#pragma acc exit data copyout(m[0:1])
+
+  for (i = 0; i < N; i++)
+    {
+      if (m->a[i] != 99)
+	abort ();
+      if (m->b[i] != 99)
+	abort ();
+      if (i > 5 && i < N-5)
+	{
+	  if (m->c[i] != 99)
+	    abort ();
+	}
+      else
+	{
+	  if (m->c[i] != 0)
+	    abort ();
+	}
+    }
+
+  free (m->a);
+  free (m->b);
+  free (m->c);
+  free (m);
+
+  return 0;
+}

libgomp/testsuite/libgomp.oacc-c++/deep-copy-13.C

@@ -0,0 +1,72 @@
+#include <stdlib.h>
+
+/* Test array slice with reference to pointer.  */
+
+typedef struct {
+  int *a;
+  int *b;
+  int *c;
+} mystruct;
+
+int main(int argc, char* argv[])
+{
+  const int N = 1024;
+  mystruct *m = (mystruct *) malloc (sizeof (*m));
+  int i;
+
+  m->a = (int *) malloc (N * sizeof (int));
+  m->b = (int *) malloc (N * sizeof (int));
+  m->c = (int *) malloc (N * sizeof (int));
+
+  for (i = 0; i < N; i++)
+    {
+      m->a[i] = 0;
+      m->b[i] = 0;
+      m->c[i] = 0;
+    }
+
+#pragma acc enter data copyin(m[0:1])
+
+  for (int i = 0; i < 99; i++)
+    {
+      int j;
+      int *&ptr = m->a;
+#pragma acc parallel loop copy(ptr[0:N])
+      for (j = 0; j < N; j++)
+	ptr[j]++;
+#pragma acc parallel loop copy(m->b[0:N], m->c[5:N-10])
+      for (j = 0; j < N; j++)
+	{
+	  m->b[j]++;
+	  if (j > 5 && j < N - 5)
+	    m->c[j]++;
+	}
+    }
+
+#pragma acc exit data copyout(m[0:1])
+
+  for (i = 0; i < N; i++)
+    {
+      if (m->a[i] != 99)
+	abort ();
+      if (m->b[i] != 99)
+	abort ();
+      if (i > 5 && i < N-5)
+	{
+	  if (m->c[i] != 99)
+	    abort ();
+	}
+      else
+	{
+	  if (m->c[i] != 0)
+	    abort ();
+	}
+    }
+
+  free (m->a);
+  free (m->b);
+  free (m->c);
+  free (m);
+
+  return 0;
+}

libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-1.c

@@ -0,0 +1,24 @@
+#include <stdlib.h>
+#include <assert.h>
+
+struct dc
+{
+  int a;
+  int *b;
+};
+
+int
+main ()
+{
+  int n = 100, i;
+  struct dc v = { .a = 3, .b = (int *) malloc (sizeof (int) * n) };
+
+#pragma acc parallel loop copy(v.a, v.b[:n])
+  for (i = 0; i < n; i++)
+    v.b[i] = v.a;
+
+  for (i = 0; i < 10; i++)
+    assert (v.b[i] == v.a);
+
+  return 0;
+}

libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c

@@ -0,0 +1,53 @@
+#include <stdlib.h>
+
+/* Test asyncronous attach and detach operation.  */
+
+typedef struct {
+  int *a;
+  int *b;
+} mystruct;
+
+int
+main (int argc, char* argv[])
+{
+  const int N = 1024;
+  mystruct m;
+  int i;
+
+  m.a = (int *) malloc (N * sizeof (int));
+  m.b = (int *) malloc (N * sizeof (int));
+
+  for (i = 0; i < N; i++)
+    {
+      m.a[i] = 0;
+      m.b[i] = 0;
+    }
+
+#pragma acc enter data copyin(m)
+
+  for (int i = 0; i < 99; i++)
+    {
+      int j;
+#pragma acc parallel loop copy(m.a[0:N]) async(i % 2)
+      for (j = 0; j < N; j++)
+	m.a[j]++;
+#pragma acc parallel loop copy(m.b[0:N]) async((i + 1) % 2)
+      for (j = 0; j < N; j++)
+	m.b[j]++;
+    }
+
+#pragma acc exit data copyout(m) wait(0, 1)
+
+  for (i = 0; i < N; i++)
+    {
+      if (m.a[i] != 99)
+	abort ();
+      if (m.b[i] != 99)
+	abort ();
+    }
+
+  free (m.a);
+  free (m.b);
+
+  return 0;
+}

libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-11.c

@@ -0,0 +1,72 @@
+#include <stdlib.h>
+
+/* Test multiple struct dereferences on one directive, and slices starting at
+   non-zero.  */
+
+typedef struct {
+  int *a;
+  int *b;
+  int *c;
+} mystruct;
+
+int main(int argc, char* argv[])
+{
+  const int N = 1024;
+  mystruct *m = (mystruct *) malloc (sizeof (*m));
+  int i;
+
+  m->a = (int *) malloc (N * sizeof (int));
+  m->b = (int *) malloc (N * sizeof (int));
+  m->c = (int *) malloc (N * sizeof (int));
+
+  for (i = 0; i < N; i++)
+    {
+      m->a[i] = 0;
+      m->b[i] = 0;
+      m->c[i] = 0;
+    }
+
+#pragma acc enter data copyin(m[0:1])
+
+  for (int i = 0; i < 99; i++)
+    {
+      int j;
+#pragma acc parallel loop copy(m->a[0:N])
+      for (j = 0; j < N; j++)
+	m->a[j]++;
+#pragma acc parallel loop copy(m->b[0:N], m->c[5:N-10])
+      for (j = 0; j < N; j++)
+	{
+	  m->b[j]++;
+	  if (j > 5 && j < N - 5)
+	    m->c[j]++;
+	}
+    }
+
+#pragma acc exit data copyout(m[0:1])
+
+  for (i = 0; i < N; i++)
+    {
+      if (m->a[i] != 99)
+	abort ();
+      if (m->b[i] != 99)
+	abort ();
+      if (i > 5 && i < N-5)
+	{
+	  if (m->c[i] != 99)
+	    abort ();
+	}
+      else
+	{
+	  if (m->c[i] != 0)
+	    abort ();
+	}
+    }
+
+  free (m->a);
+  free (m->b);
+  free (m->c);
+  free (m);
+
+  return 0;
+}

libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-14.c

@@ -0,0 +1,63 @@
+#include <openacc.h>
+#include <stdlib.h>
+
+/* Test attach/detach operation with chained dereferences.  */
+
+typedef struct mystruct {
+  int *a;
+  struct mystruct *next;
+} mystruct;
+
+int
+main (int argc, char* argv[])
+{
+  const int N = 1024;
+  mystruct *m = (mystruct *) malloc (sizeof (*m));
+  int i;
+
+  m->a = (int *) malloc (N * sizeof (int));
+  m->next = (mystruct *) malloc (sizeof (*m));
+  m->next->a = (int *) malloc (N * sizeof (int));
+  m->next->next = NULL;
+
+  for (i = 0; i < N; i++)
+    {
+      m->a[i] = 0;
+      m->next->a[i] = 0;
+    }
+
+#pragma acc enter data copyin(m[0:1])
+  acc_copyin (m->next, sizeof (*m));
+
+  for (int i = 0; i < 99; i++)
+    {
+      int j;
+      acc_copyin (m->next->a, N * sizeof (int));
+      acc_attach ((void **) &m->next);
+      /* This will attach only the innermost pointer, i.e. "a[0:N]".  That's
+	 why we have to attach the "m->next" pointer manually above.  */
+#pragma acc parallel loop copy(m->next->a[0:N])
+      for (j = 0; j < N; j++)
+	m->next->a[j]++;
+      acc_detach ((void **) &m->next);
+      acc_copyout (m->next->a, N * sizeof (int));
+    }
+
+  acc_copyout (m->next, sizeof (*m));
+#pragma acc exit data copyout(m[0:1])
+
+  for (i = 0; i < N; i++)
+    {
+      if (m->a[i] != 0)
+	abort ();
+      if (m->next->a[i] != 99)
+	abort ();
+    }
+
+  free (m->next->a);
+  free (m->next);
+  free (m->a);
+  free (m);
+
+  return 0;
+}

libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-2.c

@@ -0,0 +1,29 @@
+#include <assert.h>
+#include <stdlib.h>
+
+int
+main(int argc, char* argv[])
+{
+  struct foo {
+    int *a, *b, c, d, *e;
+  } s;
+
+  s.a = (int *) malloc (16 * sizeof (int));
+  s.b = (int *) malloc (16 * sizeof (int));
+  s.e = (int *) malloc (16 * sizeof (int));
+
+  #pragma acc data copy(s)
+  {
+    #pragma acc data copy(s.a[0:10])
+    {
+      #pragma acc parallel loop attach(s.a)
+      for (int i = 0; i < 10; i++)
+	s.a[i] = i;
+    }
+  }
+
+  for (int i = 0; i < 10; i++)
+    assert (s.a[i] == i);
+
+  return 0;
+}

libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-4.c

@@ -0,0 +1,87 @@
+#include <assert.h>
+#include <stdlib.h>
+
+#define LIST_LENGTH 10
+
+struct node
+{
+  struct node *next;
+  int val;
+};
+
+int
+sum_nodes (struct node *head)
+{
+  int i = 0, sum = 0;
+
+#pragma acc parallel reduction(+:sum) present(head[:1])
+  {
+    for (; head != NULL; head = head->next)
+      sum += head->val;
+  }
+
+  return sum;
+}
+
+void
+insert (struct node *head, int val)
+{
+  struct node *n = (struct node *) malloc (sizeof (struct node));
+
+  if (head->next)
+    {
+#pragma acc exit data detach(head->next)
+    }
+
+  n->val = val;
+  n->next = head->next;
+  head->next = n;
+
+#pragma acc enter data copyin(n[:1])
+#pragma acc enter data attach(head->next)
+  if (n->next)
+    {
+#pragma acc enter data attach(n->next)
+    }
+}
+
+void
+destroy (struct node *head)
+{
+  while (head->next != NULL)
+    {
+#pragma acc exit data detach(head->next)
+      struct node * n = head->next;
+      head->next = n->next;
+      if (n->next)
+	{
+#pragma acc exit data detach(n->next)
+	}
+#pragma acc exit data delete (n[:1])
+      if (head->next)
+	{
+#pragma acc enter data attach(head->next)
+	}
+      free (n);
+    }
+}
+
+int
+main ()
+{
+  struct node list = { .next = NULL, .val = 0 };
+  int i;
+
+#pragma acc enter data copyin(list)
+
+  for (i = 0; i < LIST_LENGTH; i++)
+    insert (&list, i + 1);
+
+  assert (sum_nodes (&list) == (LIST_LENGTH * LIST_LENGTH + LIST_LENGTH) / 2);
+
+  destroy (&list);
+
+#pragma acc exit data delete(list)
+
+  return 0;
+}

libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-6.c

@@ -0,0 +1,59 @@
+/* { dg-do run { target { ! openacc_host_selected } } } */
+
+#include <stdlib.h>
+#include <assert.h>
+#include <openacc.h>
+
+struct dc
+{
+  int a;
+  int **b;
+};
+
+int
+main ()
+{
+  int n = 100, i, j, k;
+  struct dc v = { .a = 3 };
+
+  v.b = (int **) malloc (sizeof (int *) * n);
+  for (i = 0; i < n; i++)
+    v.b[i] = (int *) malloc (sizeof (int) * n);
+
+  for (k = 0; k < 16; k++)
+    {
+#pragma acc data copy(v)
+      {
+#pragma acc data copy(v.b[:n])
+	{
+	  for (i = 0; i < n; i++)
+	    {
+	      acc_copyin (v.b[i], sizeof (int) * n);
+	      acc_attach ((void **) &v.b[i]);
+	    }
+
+#pragma acc parallel loop
+	  for (i = 0; i < n; i++)
+	    for (j = 0; j < n; j++)
+	      v.b[i][j] = v.a + i + j;
+
+	  for (i = 0; i < n; i++)
+	    {
+	      acc_detach ((void **) &v.b[i]);
+	      acc_copyout (v.b[i], sizeof (int) * n);
+	    }
+	}
+      }
+
+      for (i = 0; i < n; i++)
+	for (j = 0; j < n; j++)
+	  assert (v.b[i][j] == v.a + i + j);
+
+      assert (!acc_is_present (&v, sizeof (v)));
+      assert (!acc_is_present (v.b, sizeof (int *) * n));
+      for (i = 0; i < n; i++)
+	assert (!acc_is_present (v.b[i], sizeof (int) * n));
+    }
+
+  return 0;
+}

libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c

@@ -0,0 +1,45 @@
+/* { dg-do run { target { ! openacc_host_selected } } } */
+
+#include <stdlib.h>
+#include <assert.h>
+#include <openacc.h>
+
+struct dc
+{
+  int a;
+  int *b;
+};
+
+int
+main ()
+{
+  int n = 100, i, j, k;
+  struct dc v = { .a = 3 };
+
+  v.b = (int *) malloc (sizeof (int) * n);
+
+  for (k = 0; k < 16; k++)
+    {
+      /* Here, we do not explicitly copy the enclosing structure, but work
+	 with fields directly.  Make sure attachment counters and reference
+	 counters work properly in that case.  */
+#pragma acc enter data copyin(v.a, v.b[0:n])
+#pragma acc enter data pcopyin(v.b[0:n])
+#pragma acc enter data pcopyin(v.b[0:n])
+
+#pragma acc parallel loop present(v.a, v.b)
+      for (i = 0; i < n; i++)
+	v.b[i] = v.a + i;
+
+#pragma acc exit data copyout(v.b[:n]) finalize
+#pragma acc exit data delete(v.a)
+
+      for (i = 0; i < n; i++)
+	assert (v.b[i] == v.a + i);
+
+      assert (!acc_is_present (&v, sizeof (v)));
+      assert (!acc_is_present (v.b, sizeof (int *) * n));
+    }
+
+  return 0;
+}

libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c

@@ -0,0 +1,54 @@
+/* { dg-do run { target { ! openacc_host_selected } } } */
+
+#include <stdlib.h>
+#include <assert.h>
+#include <openacc.h>
+
+struct dc
+{
+  int a;
+  int *b;
+  int *c;
+  int *d;
+};
+
+int
+main ()
+{
+  int n = 100, i, j, k;
+  struct dc v = { .a = 3 };
+
+  v.b = (int *) malloc (sizeof (int) * n);
+  v.c = (int *) malloc (sizeof (int) * n);
+  v.d = (int *) malloc (sizeof (int) * n);
+
+#pragma acc enter data copyin(v)
+
+  for (k = 0; k < 16; k++)
+    {
+#pragma acc enter data copyin(v.a, v.b[:n], v.c[:n], v.d[:n])
+
+#pragma acc parallel loop
+      for (i = 0; i < n; i++)
+	v.b[i] = v.a + i;
+
+#pragma acc exit data copyout(v.b[:n])
+#pragma acc exit data copyout(v.c[:n])
+#pragma acc exit data copyout(v.d[:n])
+#pragma acc exit data copyout(v.a)
+
+      for (i = 0; i < n; i++)
+	assert (v.b[i] == v.a + i);
+
+      assert (acc_is_present (&v, sizeof (v)));
+      assert (!acc_is_present (v.b, sizeof (int *) * n));
+      assert (!acc_is_present (v.c, sizeof (int *) * n));
+      assert (!acc_is_present (v.d, sizeof (int *) * n));
+    }
+
+#pragma acc exit data copyout(v)
+
+  assert (!acc_is_present (&v, sizeof (v)));
+
+  return 0;
+}

libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-9.c

@@ -0,0 +1,53 @@
+#include <stdlib.h>
+
+typedef struct {
+  int *a;
+  int *b;
+} mystruct;
+
+int
+main (int argc, char* argv[])
+{
+  const int N = 1024;
+  mystruct *m = (mystruct *) malloc (sizeof (*m));
+  int i;
+
+  m->a = (int *) malloc (N * sizeof (int));
+  m->b = (int *) malloc (N * sizeof (int));
+
+  for (i = 0; i < N; i++)
+    {
+      m->a[i] = 0;
+      m->b[i] = 0;
+    }
+
+#pragma acc enter data copyin(m[0:1])
+
+  for (int i = 0; i < 99; i++)
+    {
+      int j;
+      int *ptr = m->a;
+#pragma acc parallel loop copy(m->a[0:N])
+      for (j = 0; j < N; j++)
+	m->a[j]++;
+#pragma acc parallel loop copy(m->b[0:N])
+      for (j = 0; j < N; j++)
+	m->b[j]++;
+    }
+
+#pragma acc exit data copyout(m[0:1])
+
+  for (i = 0; i < N; i++)
+    {
+      if (m->a[i] != 99)
+	abort ();
+      if (m->b[i] != 99)
+	abort ();
+    }
+
+  free (m->a);
+  free (m->b);
+  free (m);
+
+  return 0;
+}