Skip to content

Commit c71da7d

Browse files
authored
[OpenMP] Add tests for mapping of chained 'containing' structs (#156703)
This PR adds several new tests for mapping of chained structures, i.e. those resembling: #pragma omp target map(tofrom: a->b->c) These are currently XFAILed, although the first two tests actually work with unified memory -- I'm not sure if it's possible to easily improve the condition on the XFAILs in question to make them more accurate. These cases are all fixed by the WIP PR #153683.
1 parent c173476 commit c71da7d

File tree

3 files changed

+351
-0
lines changed

3 files changed

+351
-0
lines changed
Lines changed: 58 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,58 @@
1+
// RUN: %libomptarget-compilexx-run-and-check-generic
2+
// XFAIL: *
3+
4+
#include <cstdlib>
5+
#include <cstdio>
6+
#include <cassert>
7+
8+
struct S {
9+
int a;
10+
int b;
11+
int c;
12+
};
13+
14+
struct T {
15+
S *s0;
16+
S *s1;
17+
S *s2;
18+
};
19+
20+
int main() {
21+
T *v = (T *) malloc (sizeof(T));
22+
v->s0 = (S *) malloc (sizeof(S));
23+
v->s1 = (S *) malloc (sizeof(S));
24+
v->s2 = (S *) malloc (sizeof(S));
25+
v->s0->a = 10;
26+
v->s0->b = 10;
27+
v->s0->c = 10;
28+
v->s1->a = 20;
29+
v->s1->b = 20;
30+
v->s1->c = 20;
31+
v->s2->a = 30;
32+
v->s2->b = 30;
33+
v->s2->c = 30;
34+
35+
#pragma omp target map(to: v[:1]) map(tofrom: v->s1->b, v->s1->c, v->s2->b)
36+
{
37+
v->s1->b += 3;
38+
v->s1->c += 5;
39+
v->s2->b += 7;
40+
}
41+
42+
printf ("%d\n", v->s0->a); // CHECK: 10
43+
printf ("%d\n", v->s0->b); // CHECK: 10
44+
printf ("%d\n", v->s0->c); // CHECK: 10
45+
printf ("%d\n", v->s1->a); // CHECK: 20
46+
printf ("%d\n", v->s1->b); // CHECK: 23
47+
printf ("%d\n", v->s1->c); // CHECK: 25
48+
printf ("%d\n", v->s2->a); // CHECK: 30
49+
printf ("%d\n", v->s2->b); // CHECK: 37
50+
printf ("%d\n", v->s2->c); // CHECK: 30
51+
52+
free(v->s0);
53+
free(v->s1);
54+
free(v->s2);
55+
free(v);
56+
57+
return 0;
58+
}
Lines changed: 76 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,76 @@
1+
// RUN: %libomptarget-compilexx-run-and-check-generic
2+
// XFAIL: *
3+
4+
#include <cstdlib>
5+
#include <cstdio>
6+
#include <cassert>
7+
8+
struct R {
9+
int d;
10+
int e;
11+
int f;
12+
};
13+
14+
struct S {
15+
R *r0;
16+
R *r1;
17+
R *r2;
18+
};
19+
20+
struct T {
21+
S *s0;
22+
S *s1;
23+
S *s2;
24+
};
25+
26+
int main() {
27+
T *v = (T *) malloc (sizeof(T));
28+
29+
v->s0 = (S *) malloc (sizeof(S));
30+
v->s1 = (S *) malloc (sizeof(S));
31+
v->s2 = (S *) malloc (sizeof(S));
32+
33+
v->s0->r0 = (R *) calloc (1, sizeof(R));
34+
v->s0->r1 = (R *) calloc (1, sizeof(R));
35+
v->s0->r2 = (R *) calloc (1, sizeof(R));
36+
37+
v->s1->r0 = (R *) calloc (1, sizeof(R));
38+
v->s1->r1 = (R *) calloc (1, sizeof(R));
39+
v->s1->r2 = (R *) calloc (1, sizeof(R));
40+
41+
v->s2->r0 = (R *) calloc (1, sizeof(R));
42+
v->s2->r1 = (R *) calloc (1, sizeof(R));
43+
v->s2->r2 = (R *) calloc (1, sizeof(R));
44+
45+
#pragma omp target map(to: v->s1, v->s2, *v->s1, v->s1->r1, *v->s2, v->s2->r0) \
46+
map(tofrom: v->s1->r1->d, v->s1->r1->e, v->s1->r2->d, v->s1->r2->f, v->s2->r0->e)
47+
{
48+
v->s1->r1->d += 3;
49+
v->s1->r1->e += 5;
50+
v->s1->r2->d += 7;
51+
v->s1->r2->f += 9;
52+
v->s2->r0->e += 11;
53+
}
54+
55+
printf ("%d\n", v->s1->r1->d); // CHECK: 3
56+
printf ("%d\n", v->s1->r1->e); // CHECK: 5
57+
printf ("%d\n", v->s1->r2->d); // CHECK: 7
58+
printf ("%d\n", v->s1->r2->f); // CHECK: 9
59+
printf ("%d\n", v->s2->r0->e); // CHECK: 11
60+
61+
free(v->s0->r0);
62+
free(v->s0->r1);
63+
free(v->s0->r2);
64+
free(v->s1->r0);
65+
free(v->s1->r1);
66+
free(v->s1->r2);
67+
free(v->s2->r0);
68+
free(v->s2->r1);
69+
free(v->s2->r2);
70+
free(v->s0);
71+
free(v->s1);
72+
free(v->s2);
73+
free(v);
74+
75+
return 0;
76+
}
Lines changed: 217 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,217 @@
1+
// RUN: %libomptarget-compilexx-run-and-check-generic
2+
3+
// XFAIL: *
4+
5+
#include <cstdlib>
6+
#include <cstdio>
7+
#include <cassert>
8+
#include <cstring>
9+
10+
#include <omp.h>
11+
12+
struct R {
13+
int d;
14+
int e;
15+
int f;
16+
};
17+
18+
struct S {
19+
int a;
20+
int b;
21+
struct {
22+
int c;
23+
R r;
24+
R *rp;
25+
} sub;
26+
int g;
27+
};
28+
29+
struct T {
30+
int a;
31+
int *ptr;
32+
int b;
33+
};
34+
35+
int main() {
36+
R r;
37+
R *rp = new R;
38+
S s;
39+
S *sp = new S;
40+
T t;
41+
T *tp = new T;
42+
43+
memset(&r, 0, sizeof(R));
44+
memset(rp, 0, sizeof(R));
45+
memset(&s, 0, sizeof(S));
46+
memset(sp, 0, sizeof(S));
47+
memset(&t, 0, sizeof(T));
48+
memset(tp, 0, sizeof(T));
49+
50+
s.sub.rp = new R;
51+
sp->sub.rp = new R;
52+
53+
memset(s.sub.rp, 0, sizeof(R));
54+
memset(sp->sub.rp, 0, sizeof(R));
55+
56+
t.ptr = new int[10];
57+
tp->ptr = new int[10];
58+
59+
memset(t.ptr, 0, sizeof(int)*10);
60+
memset(tp->ptr, 0, sizeof(int)*10);
61+
62+
#pragma omp target map(tofrom: r) map(tofrom: r.e)
63+
{
64+
r.d++;
65+
r.e += 2;
66+
r.f += 3;
67+
}
68+
printf ("%d\n", r.d); // CHECK: 1
69+
printf ("%d\n", r.e); // CHECK-NEXT: 2
70+
printf ("%d\n", r.f); // CHECK-NEXT: 3
71+
72+
#pragma omp target map(tofrom: rp[:1]) map(tofrom: rp->e)
73+
{
74+
rp->d++;
75+
rp->e += 2;
76+
rp->f += 3;
77+
}
78+
79+
printf ("%d\n", rp->d); // CHECK-NEXT: 1
80+
printf ("%d\n", rp->e); // CHECK-NEXT: 2
81+
printf ("%d\n", rp->f); // CHECK-NEXT: 3
82+
83+
int v;
84+
int *orig_addr_v = &v;
85+
bool separate_memory_space;
86+
87+
#pragma omp target data map(v)
88+
{
89+
void *mapped_ptr_v =
90+
omp_get_mapped_ptr(orig_addr_v, omp_get_default_device());
91+
separate_memory_space = mapped_ptr_v != (void*) orig_addr_v;
92+
}
93+
94+
const char *mapping_flavour = separate_memory_space ? "separate" : "unified";
95+
96+
#pragma omp target map(to: s) map(tofrom: s.sub.r.e)
97+
{
98+
s.b++;
99+
s.sub.r.d+=2;
100+
s.sub.r.e+=3;
101+
s.sub.r.f+=4;
102+
}
103+
104+
printf ("%d/%s\n", s.b, mapping_flavour);
105+
printf ("%d/%s\n", s.sub.r.d, mapping_flavour);
106+
printf ("%d/%s\n", s.sub.r.e, mapping_flavour);
107+
printf ("%d/%s\n", s.sub.r.f, mapping_flavour);
108+
109+
// CHECK: {{0/separate|1/unified}}
110+
// CHECK-NEXT: {{0/separate|2/unified}}
111+
// CHECK-NEXT: 3
112+
// CHECK-NEXT: {{0/separate|4/unified}}
113+
114+
#pragma omp target map(to: s, s.b) map(to: s.sub.rp[:1]) map(tofrom: s.sub.rp->e)
115+
{
116+
s.b++;
117+
s.sub.rp->d+=2;
118+
s.sub.rp->e+=3;
119+
s.sub.rp->f+=4;
120+
}
121+
122+
printf ("%d/%s\n", s.b, mapping_flavour);
123+
printf ("%d/%s\n", s.sub.rp->d, mapping_flavour);
124+
printf ("%d/%s\n", s.sub.rp->e, mapping_flavour);
125+
printf ("%d/%s\n", s.sub.rp->f, mapping_flavour);
126+
127+
// CHECK-NEXT: {{0/separate|2/unified}}
128+
// CHECK-NEXT: {{0/separate|2/unified}}
129+
// CHECK-NEXT: 3
130+
// CHECK-NEXT: {{0/separate|4/unified}}
131+
132+
#pragma omp target map(to: sp[:1]) map(tofrom: sp->sub.r.e)
133+
{
134+
sp->b++;
135+
sp->sub.r.d+=2;
136+
sp->sub.r.e+=3;
137+
sp->sub.r.f+=4;
138+
}
139+
140+
printf ("%d/%s\n", sp->b, mapping_flavour);
141+
printf ("%d/%s\n", sp->sub.r.d, mapping_flavour);
142+
printf ("%d/%s\n", sp->sub.r.e, mapping_flavour);
143+
printf ("%d/%s\n", sp->sub.r.f, mapping_flavour);
144+
145+
// CHECK-NEXT: {{0/separate|1/unified}}
146+
// CHECK-NEXT: {{0/separate|2/unified}}
147+
// CHECK-NEXT: 3
148+
// CHECK-NEXT: {{0/separate|4/unified}}
149+
150+
#pragma omp target map(to: sp[:1]) map(to: sp->sub.rp[:1]) map(tofrom: sp->sub.rp->e)
151+
{
152+
sp->b++;
153+
sp->sub.rp->d+=2;
154+
sp->sub.rp->e+=3;
155+
sp->sub.rp->f+=4;
156+
}
157+
158+
printf ("%d/%s\n", sp->b, mapping_flavour);
159+
printf ("%d/%s\n", sp->sub.rp->d, mapping_flavour);
160+
printf ("%d/%s\n", sp->sub.rp->e, mapping_flavour);
161+
printf ("%d/%s\n", sp->sub.rp->f, mapping_flavour);
162+
163+
// CHECK-NEXT: {{0/separate|2/unified}}
164+
// CHECK-NEXT: {{0/separate|2/unified}}
165+
// CHECK-NEXT: 3
166+
// CHECK-NEXT: {{0/separate|4/unified}}
167+
168+
#pragma omp target map(tofrom: t) map(tofrom: t.ptr[2:1])
169+
{
170+
t.a++;
171+
t.ptr[2]+=2;
172+
t.b+=3;
173+
}
174+
175+
printf ("%d\n", t.a); // CHECK-NEXT: 1
176+
printf ("%d\n", t.ptr[2]); // CHECK-NEXT: 2
177+
printf ("%d\n", t.b); // CHECK-NEXT: 3
178+
179+
#pragma omp target map(tofrom: t) map(tofrom: t.a)
180+
{
181+
t.b++;
182+
}
183+
184+
printf ("%d\n", t.b); // CHECK-NEXT: 4
185+
186+
#pragma omp target map(tofrom: t) map(tofrom: t.ptr[2:1], t.a)
187+
{
188+
t.a++;
189+
t.ptr[2]+=2;
190+
t.b+=3;
191+
}
192+
193+
printf ("%d\n", t.a); // CHECK-NEXT: 2
194+
printf ("%d\n", t.ptr[2]); // CHECK-NEXT: 4
195+
printf ("%d\n", t.b); // CHECK-NEXT: 7
196+
197+
#pragma omp target map(tofrom: t) map(tofrom: t.ptr[2:1], t.a)
198+
{
199+
/* Empty */
200+
}
201+
202+
printf ("%d\n", t.a); // CHECK-NEXT: 2
203+
printf ("%d\n", t.ptr[2]); // CHECK-NEXT: 4
204+
printf ("%d\n", t.b); // CHECK-NEXT: 7
205+
206+
delete s.sub.rp;
207+
delete sp->sub.rp;
208+
209+
delete[] t.ptr;
210+
delete[] tp->ptr;
211+
212+
delete rp;
213+
delete sp;
214+
delete tp;
215+
216+
return 0;
217+
}

0 commit comments

Comments
 (0)