1 #include <stdlib.h>
2
3 /* Test mapping chained indirect struct accesses, mixed in different ways. */
4
5 typedef struct {
6 int *a;
7 int b;
8 int *c;
9 } str1;
10
11 typedef struct {
12 int d;
13 int *e;
14 str1 *f;
15 } str2;
16
17 typedef struct {
18 int g;
19 int h;
20 str2 *s2;
21 } str3;
22
23 typedef struct {
24 str3 m;
25 str3 n;
26 } str4;
27
28 void
29 zero_arrays (str4 *s, int N)
30 {
31 for (int i = 0; i < N; i++)
32 {
33 s->m.s2->e[i] = 0;
34 s->m.s2->f->a[i] = 0;
35 s->m.s2->f->c[i] = 0;
36 s->n.s2->e[i] = 0;
37 s->n.s2->f->a[i] = 0;
38 s->n.s2->f->c[i] = 0;
39 }
40 }
41
42 void
43 alloc_s2 (str2 **s, int N)
44 {
45 (*s) = (str2 *) malloc (sizeof (str2));
46 (*s)->f = (str1 *) malloc (sizeof (str1));
47 (*s)->e = (int *) malloc (sizeof (int) * N);
48 (*s)->f->a = (int *) malloc (sizeof (int) * N);
49 (*s)->f->c = (int *) malloc (sizeof (int) * N);
50 }
51
52 int main (int argc, char* argv[])
53 {
54 const int N = 1024;
55 str4 p, *q;
56 int i;
57
58 alloc_s2 (&p.m.s2, N);
59 alloc_s2 (&p.n.s2, N);
60 q = (str4 *) malloc (sizeof (str4));
61 alloc_s2 (&q->m.s2, N);
62 alloc_s2 (&q->n.s2, N);
63
64 zero_arrays (&p, N);
65
66 for (int i = 0; i < 99; i++)
67 {
68 #pragma acc enter data copyin(p.m.s2[:1])
69 #pragma acc parallel loop copy(p.m.s2->e[:N])
70 for (int j = 0; j < N; j++)
71 p.m.s2->e[j]++;
72 #pragma acc exit data delete(p.m.s2[:1])
73 }
74
75 for (i = 0; i < N; i++)
76 if (p.m.s2->e[i] != 99)
77 abort ();
78
79 zero_arrays (&p, N);
80
81 for (int i = 0; i < 99; i++)
82 {
83 #pragma acc enter data copyin(p.m.s2[:1])
84 #pragma acc enter data copyin(p.m.s2->f[:1])
85 #pragma acc parallel loop copy(p.m.s2->f->a[:N]) copy(p.m.s2->f->c[:N])
86 for (int j = 0; j < N; j++)
87 {
88 p.m.s2->f->a[j]++;
89 p.m.s2->f->c[j]++;
90 }
91 #pragma acc exit data delete(p.m.s2->f[:1])
92 #pragma acc exit data delete(p.m.s2[:1])
93 }
94
95 for (i = 0; i < N; i++)
96 if (p.m.s2->f->a[i] != 99 || p.m.s2->f->c[i] != 99)
97 abort ();
98
99 zero_arrays (&p, N);
100
101 for (int i = 0; i < 99; i++)
102 {
103 #pragma acc enter data copyin(p.m.s2[:1]) copyin(p.n.s2[:1])
104 #pragma acc enter data copyin(p.m.s2->f[:1]) copyin(p.n.s2->f[:1])
105 #pragma acc parallel loop copy(p.m.s2->f->a[:N]) copy(p.m.s2->f->c[:N]) \
106 copy(p.n.s2->f->a[:N]) copy(p.n.s2->f->c[:N])
107 for (int j = 0; j < N; j++)
108 {
109 p.m.s2->f->a[j]++;
110 p.m.s2->f->c[j]++;
111 p.n.s2->f->a[j]++;
112 p.n.s2->f->c[j]++;
113 }
114 #pragma acc exit data delete(p.m.s2->f[:1]) delete(p.n.s2->f[:1])
115 #pragma acc exit data delete(p.m.s2[:1]) delete(p.n.s2[:1])
116 }
117
118 for (i = 0; i < N; i++)
119 if (p.m.s2->f->a[i] != 99 || p.m.s2->f->c[i] != 99
120 || p.n.s2->f->a[i] != 99 || p.n.s2->f->c[i] != 99)
121 abort ();
122
123 zero_arrays (&p, N);
124
125 for (int i = 0; i < 99; i++)
126 {
127 #pragma acc enter data copyin(p.m.s2[:1]) copyin(p.n.s2[:1])
128 #pragma acc enter data copyin(p.n.s2->e[:N]) copyin(p.n.s2->f[:1]) \
129 copyin(p.m.s2->f[:1])
130 #pragma acc parallel loop copy(p.m.s2->f->a[:N]) copy(p.n.s2->f->a[:N])
131 for (int j = 0; j < N; j++)
132 {
133 p.m.s2->f->a[j]++;
134 p.n.s2->f->a[j]++;
135 p.n.s2->e[j]++;
136 }
137 #pragma acc exit data delete(p.m.s2->f[:1]) delete(p.n.s2->f[:1]) \
138 copyout(p.n.s2->e[:N])
139 #pragma acc exit data delete(p.m.s2[:1]) delete(p.n.s2[:1])
140 }
141
142 for (i = 0; i < N; i++)
143 if (p.m.s2->f->a[i] != 99 || p.n.s2->f->a[i] != 99
144 || p.n.s2->e[i] != 99)
145 abort ();
146
147 zero_arrays (q, N);
148
149 for (int i = 0; i < 99; i++)
150 {
151 #pragma acc enter data copyin(q->m.s2[:1])
152 #pragma acc parallel loop copy(q->m.s2->e[:N])
153 for (int j = 0; j < N; j++)
154 q->m.s2->e[j]++;
155 #pragma acc exit data delete(q->m.s2[:1])
156 }
157
158 for (i = 0; i < N; i++)
159 if (q->m.s2->e[i] != 99)
160 abort ();
161
162 zero_arrays (q, N);
163
164 for (int i = 0; i < 99; i++)
165 {
166 #pragma acc enter data copyin(q->m.s2[:1])
167 #pragma acc enter data copyin(q->m.s2->f[:1])
168 #pragma acc parallel loop copy(q->m.s2->f->a[:N]) copy(q->m.s2->f->c[:N])
169 for (int j = 0; j < N; j++)
170 {
171 q->m.s2->f->a[j]++;
172 q->m.s2->f->c[j]++;
173 }
174 #pragma acc exit data delete(q->m.s2->f[:1])
175 #pragma acc exit data delete(q->m.s2[:1])
176 }
177
178 for (i = 0; i < N; i++)
179 if (q->m.s2->f->a[i] != 99 || q->m.s2->f->c[i] != 99)
180 abort ();
181
182 zero_arrays (q, N);
183
184 for (int i = 0; i < 99; i++)
185 {
186 #pragma acc enter data copyin(q->m.s2[:1]) copyin(q->n.s2[:1])
187 #pragma acc enter data copyin(q->m.s2->f[:1]) copyin(q->n.s2->f[:1])
188 #pragma acc parallel loop copy(q->m.s2->f->a[:N]) copy(q->m.s2->f->c[:N]) \
189 copy(q->n.s2->f->a[:N]) copy(q->n.s2->f->c[:N])
190 for (int j = 0; j < N; j++)
191 {
192 q->m.s2->f->a[j]++;
193 q->m.s2->f->c[j]++;
194 q->n.s2->f->a[j]++;
195 q->n.s2->f->c[j]++;
196 }
197 #pragma acc exit data delete(q->m.s2->f[:1]) delete(q->n.s2->f[:1])
198 #pragma acc exit data delete(q->m.s2[:1]) delete(q->n.s2[:1])
199 }
200
201 for (i = 0; i < N; i++)
202 if (q->m.s2->f->a[i] != 99 || q->m.s2->f->c[i] != 99
203 || q->n.s2->f->a[i] != 99 || q->n.s2->f->c[i] != 99)
204 abort ();
205
206 zero_arrays (q, N);
207
208 for (int i = 0; i < 99; i++)
209 {
210 #pragma acc enter data copyin(q->m.s2[:1]) copyin(q->n.s2[:1])
211 #pragma acc enter data copyin(q->n.s2->e[:N]) copyin(q->m.s2->f[:1]) \
212 copyin(q->n.s2->f[:1])
213 #pragma acc parallel loop copy(q->m.s2->f->a[:N]) copy(q->n.s2->f->a[:N])
214 for (int j = 0; j < N; j++)
215 {
216 q->m.s2->f->a[j]++;
217 q->n.s2->f->a[j]++;
218 q->n.s2->e[j]++;
219 }
220 #pragma acc exit data delete(q->m.s2->f[:1]) delete(q->n.s2->f[:1]) \
221 copyout(q->n.s2->e[:N])
222 #pragma acc exit data delete(q->m.s2[:1]) delete(q->n.s2[:1])
223 }
224
225 for (i = 0; i < N; i++)
226 if (q->m.s2->f->a[i] != 99 || q->n.s2->f->a[i] != 99
227 || q->n.s2->e[i] != 99)
228 abort ();
229
230 return 0;
231 }