This patch has been broken out of the "OpenACC 2.6 manual deep copy support" patch, last posted here:
https://gcc.gnu.org/ml/gcc-patches/2019-11/msg02376.html This part adds C and C++ execution tests to libgomp. Tested alongside other patches in this series with offloading to NVPTX. OK? Thanks, Julian ChangeLog 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. --- .../testsuite/libgomp.oacc-c++/deep-copy-12.C | 72 +++++++++++++++ .../testsuite/libgomp.oacc-c++/deep-copy-13.C | 72 +++++++++++++++ .../libgomp.oacc-c-c++-common/deep-copy-1.c | 24 +++++ .../libgomp.oacc-c-c++-common/deep-copy-10.c | 53 +++++++++++ .../libgomp.oacc-c-c++-common/deep-copy-11.c | 72 +++++++++++++++ .../libgomp.oacc-c-c++-common/deep-copy-14.c | 63 ++++++++++++++ .../libgomp.oacc-c-c++-common/deep-copy-2.c | 29 +++++++ .../libgomp.oacc-c-c++-common/deep-copy-4.c | 87 +++++++++++++++++++ .../libgomp.oacc-c-c++-common/deep-copy-6.c | 59 +++++++++++++ .../libgomp.oacc-c-c++-common/deep-copy-7.c | 45 ++++++++++ .../libgomp.oacc-c-c++-common/deep-copy-8.c | 54 ++++++++++++ .../libgomp.oacc-c-c++-common/deep-copy-9.c | 53 +++++++++++ 12 files changed, 683 insertions(+) create mode 100644 libgomp/testsuite/libgomp.oacc-c++/deep-copy-12.C create mode 100644 libgomp/testsuite/libgomp.oacc-c++/deep-copy-13.C create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-1.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-11.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-14.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-2.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-4.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-6.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-9.c diff --git a/libgomp/testsuite/libgomp.oacc-c++/deep-copy-12.C b/libgomp/testsuite/libgomp.oacc-c++/deep-copy-12.C new file mode 100644 index 00000000000..a512008685d --- /dev/null +++ b/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; +} diff --git a/libgomp/testsuite/libgomp.oacc-c++/deep-copy-13.C b/libgomp/testsuite/libgomp.oacc-c++/deep-copy-13.C new file mode 100644 index 00000000000..a5194568603 --- /dev/null +++ b/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; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-1.c new file mode 100644 index 00000000000..d8d7067e452 --- /dev/null +++ b/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; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c new file mode 100644 index 00000000000..573a8214bf0 --- /dev/null +++ b/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; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-11.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-11.c new file mode 100644 index 00000000000..db6012fb352 --- /dev/null +++ b/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; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-14.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-14.c new file mode 100644 index 00000000000..275fa9ae256 --- /dev/null +++ b/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; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-2.c new file mode 100644 index 00000000000..7e26e9aa8b9 --- /dev/null +++ b/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; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-4.c new file mode 100644 index 00000000000..8874ca0a504 --- /dev/null +++ b/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; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-6.c new file mode 100644 index 00000000000..391149459c9 --- /dev/null +++ b/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; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c new file mode 100644 index 00000000000..a59047af520 --- /dev/null +++ b/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; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c new file mode 100644 index 00000000000..0ca5990b377 --- /dev/null +++ b/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; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-9.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-9.c new file mode 100644 index 00000000000..e86a46bd84a --- /dev/null +++ b/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; +} -- 2.23.0
