| #include <stdlib.h> |
| |
| /* Test mapping chained indirect struct accesses, mixed in different ways. */ |
| |
| typedef struct { |
| int *a; |
| int b; |
| int *c; |
| } str1; |
| |
| typedef struct { |
| int d; |
| int *e; |
| str1 *f; |
| } str2; |
| |
| typedef struct { |
| int g; |
| int h; |
| str2 *s2; |
| } str3; |
| |
| typedef struct { |
| str3 m; |
| str3 n; |
| } str4; |
| |
| void |
| zero_arrays (str4 *s, int N) |
| { |
| for (int i = 0; i < N; i++) |
| { |
| s->m.s2->e[i] = 0; |
| s->m.s2->f->a[i] = 0; |
| s->m.s2->f->c[i] = 0; |
| s->n.s2->e[i] = 0; |
| s->n.s2->f->a[i] = 0; |
| s->n.s2->f->c[i] = 0; |
| } |
| } |
| |
| void |
| alloc_s2 (str2 **s, int N) |
| { |
| (*s) = (str2 *) malloc (sizeof (str2)); |
| (*s)->f = (str1 *) malloc (sizeof (str1)); |
| (*s)->e = (int *) malloc (sizeof (int) * N); |
| (*s)->f->a = (int *) malloc (sizeof (int) * N); |
| (*s)->f->c = (int *) malloc (sizeof (int) * N); |
| } |
| |
| int main (int argc, char* argv[]) |
| { |
| const int N = 1024; |
| str4 p, *q; |
| int i; |
| |
| alloc_s2 (&p.m.s2, N); |
| alloc_s2 (&p.n.s2, N); |
| q = (str4 *) malloc (sizeof (str4)); |
| alloc_s2 (&q->m.s2, N); |
| alloc_s2 (&q->n.s2, N); |
| |
| zero_arrays (&p, N); |
| |
| for (int i = 0; i < 99; i++) |
| { |
| #pragma acc enter data copyin(p.m.s2[:1]) |
| #pragma acc parallel loop copy(p.m.s2->e[:N]) |
| for (int j = 0; j < N; j++) |
| p.m.s2->e[j]++; |
| #pragma acc exit data delete(p.m.s2[:1]) |
| } |
| |
| for (i = 0; i < N; i++) |
| if (p.m.s2->e[i] != 99) |
| abort (); |
| |
| zero_arrays (&p, N); |
| |
| for (int i = 0; i < 99; i++) |
| { |
| #pragma acc enter data copyin(p.m.s2[:1]) |
| #pragma acc enter data copyin(p.m.s2->f[:1]) |
| #pragma acc parallel loop copy(p.m.s2->f->a[:N]) copy(p.m.s2->f->c[:N]) |
| for (int j = 0; j < N; j++) |
| { |
| p.m.s2->f->a[j]++; |
| p.m.s2->f->c[j]++; |
| } |
| #pragma acc exit data delete(p.m.s2->f[:1]) |
| #pragma acc exit data delete(p.m.s2[:1]) |
| } |
| |
| for (i = 0; i < N; i++) |
| if (p.m.s2->f->a[i] != 99 || p.m.s2->f->c[i] != 99) |
| abort (); |
| |
| zero_arrays (&p, N); |
| |
| for (int i = 0; i < 99; i++) |
| { |
| #pragma acc enter data copyin(p.m.s2[:1]) copyin(p.n.s2[:1]) |
| #pragma acc enter data copyin(p.m.s2->f[:1]) copyin(p.n.s2->f[:1]) |
| #pragma acc parallel loop copy(p.m.s2->f->a[:N]) copy(p.m.s2->f->c[:N]) \ |
| copy(p.n.s2->f->a[:N]) copy(p.n.s2->f->c[:N]) |
| for (int j = 0; j < N; j++) |
| { |
| p.m.s2->f->a[j]++; |
| p.m.s2->f->c[j]++; |
| p.n.s2->f->a[j]++; |
| p.n.s2->f->c[j]++; |
| } |
| #pragma acc exit data delete(p.m.s2->f[:1]) delete(p.n.s2->f[:1]) |
| #pragma acc exit data delete(p.m.s2[:1]) delete(p.n.s2[:1]) |
| } |
| |
| for (i = 0; i < N; i++) |
| if (p.m.s2->f->a[i] != 99 || p.m.s2->f->c[i] != 99 |
| || p.n.s2->f->a[i] != 99 || p.n.s2->f->c[i] != 99) |
| abort (); |
| |
| zero_arrays (&p, N); |
| |
| for (int i = 0; i < 99; i++) |
| { |
| #pragma acc enter data copyin(p.m.s2[:1]) copyin(p.n.s2[:1]) |
| #pragma acc enter data copyin(p.n.s2->e[:N]) copyin(p.n.s2->f[:1]) \ |
| copyin(p.m.s2->f[:1]) |
| #pragma acc parallel loop copy(p.m.s2->f->a[:N]) copy(p.n.s2->f->a[:N]) |
| for (int j = 0; j < N; j++) |
| { |
| p.m.s2->f->a[j]++; |
| p.n.s2->f->a[j]++; |
| p.n.s2->e[j]++; |
| } |
| #pragma acc exit data delete(p.m.s2->f[:1]) delete(p.n.s2->f[:1]) \ |
| copyout(p.n.s2->e[:N]) |
| #pragma acc exit data delete(p.m.s2[:1]) delete(p.n.s2[:1]) |
| } |
| |
| for (i = 0; i < N; i++) |
| if (p.m.s2->f->a[i] != 99 || p.n.s2->f->a[i] != 99 |
| || p.n.s2->e[i] != 99) |
| abort (); |
| |
| zero_arrays (q, N); |
| |
| for (int i = 0; i < 99; i++) |
| { |
| #pragma acc enter data copyin(q->m.s2[:1]) |
| #pragma acc parallel loop copy(q->m.s2->e[:N]) |
| for (int j = 0; j < N; j++) |
| q->m.s2->e[j]++; |
| #pragma acc exit data delete(q->m.s2[:1]) |
| } |
| |
| for (i = 0; i < N; i++) |
| if (q->m.s2->e[i] != 99) |
| abort (); |
| |
| zero_arrays (q, N); |
| |
| for (int i = 0; i < 99; i++) |
| { |
| #pragma acc enter data copyin(q->m.s2[:1]) |
| #pragma acc enter data copyin(q->m.s2->f[:1]) |
| #pragma acc parallel loop copy(q->m.s2->f->a[:N]) copy(q->m.s2->f->c[:N]) |
| for (int j = 0; j < N; j++) |
| { |
| q->m.s2->f->a[j]++; |
| q->m.s2->f->c[j]++; |
| } |
| #pragma acc exit data delete(q->m.s2->f[:1]) |
| #pragma acc exit data delete(q->m.s2[:1]) |
| } |
| |
| for (i = 0; i < N; i++) |
| if (q->m.s2->f->a[i] != 99 || q->m.s2->f->c[i] != 99) |
| abort (); |
| |
| zero_arrays (q, N); |
| |
| for (int i = 0; i < 99; i++) |
| { |
| #pragma acc enter data copyin(q->m.s2[:1]) copyin(q->n.s2[:1]) |
| #pragma acc enter data copyin(q->m.s2->f[:1]) copyin(q->n.s2->f[:1]) |
| #pragma acc parallel loop copy(q->m.s2->f->a[:N]) copy(q->m.s2->f->c[:N]) \ |
| copy(q->n.s2->f->a[:N]) copy(q->n.s2->f->c[:N]) |
| for (int j = 0; j < N; j++) |
| { |
| q->m.s2->f->a[j]++; |
| q->m.s2->f->c[j]++; |
| q->n.s2->f->a[j]++; |
| q->n.s2->f->c[j]++; |
| } |
| #pragma acc exit data delete(q->m.s2->f[:1]) delete(q->n.s2->f[:1]) |
| #pragma acc exit data delete(q->m.s2[:1]) delete(q->n.s2[:1]) |
| } |
| |
| for (i = 0; i < N; i++) |
| if (q->m.s2->f->a[i] != 99 || q->m.s2->f->c[i] != 99 |
| || q->n.s2->f->a[i] != 99 || q->n.s2->f->c[i] != 99) |
| abort (); |
| |
| zero_arrays (q, N); |
| |
| for (int i = 0; i < 99; i++) |
| { |
| #pragma acc enter data copyin(q->m.s2[:1]) copyin(q->n.s2[:1]) |
| #pragma acc enter data copyin(q->n.s2->e[:N]) copyin(q->m.s2->f[:1]) \ |
| copyin(q->n.s2->f[:1]) |
| #pragma acc parallel loop copy(q->m.s2->f->a[:N]) copy(q->n.s2->f->a[:N]) |
| for (int j = 0; j < N; j++) |
| { |
| q->m.s2->f->a[j]++; |
| q->n.s2->f->a[j]++; |
| q->n.s2->e[j]++; |
| } |
| #pragma acc exit data delete(q->m.s2->f[:1]) delete(q->n.s2->f[:1]) \ |
| copyout(q->n.s2->e[:N]) |
| #pragma acc exit data delete(q->m.s2[:1]) delete(q->n.s2[:1]) |
| } |
| |
| for (i = 0; i < N; i++) |
| if (q->m.s2->f->a[i] != 99 || q->n.s2->f->a[i] != 99 |
| || q->n.s2->e[i] != 99) |
| abort (); |
| |
| return 0; |
| } |