+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.
--- /dev/null
+#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;
+}
--- /dev/null
+#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;
+}
--- /dev/null
+#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;
+}
--- /dev/null
+#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;
+}
--- /dev/null
+#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;
+}
--- /dev/null
+#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;
+}
--- /dev/null
+#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;
+}
--- /dev/null
+#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;
+}
--- /dev/null
+/* { 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;
+}
--- /dev/null
+/* { 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;
+}
--- /dev/null
+/* { 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;
+}
--- /dev/null
+#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;
+}
git://gcc.gnu.org / gcc.git / commitdiff