diff --git a/offload/test/mapping/chained_containing_structs_1.cc b/offload/test/mapping/chained_containing_structs_1.cc new file mode 100644 index 0000000000000..4dbb17140de12 --- /dev/null +++ b/offload/test/mapping/chained_containing_structs_1.cc @@ -0,0 +1,58 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic +// XFAIL: * + +#include +#include +#include + +struct S { + int a; + int b; + int c; +}; + +struct T { + S *s0; + S *s1; + S *s2; +}; + +int main() { + T *v = (T *) malloc (sizeof(T)); + v->s0 = (S *) malloc (sizeof(S)); + v->s1 = (S *) malloc (sizeof(S)); + v->s2 = (S *) malloc (sizeof(S)); + v->s0->a = 10; + v->s0->b = 10; + v->s0->c = 10; + v->s1->a = 20; + v->s1->b = 20; + v->s1->c = 20; + v->s2->a = 30; + v->s2->b = 30; + v->s2->c = 30; + +#pragma omp target map(to: v[:1]) map(tofrom: v->s1->b, v->s1->c, v->s2->b) + { + v->s1->b += 3; + v->s1->c += 5; + v->s2->b += 7; + } + + printf ("%d\n", v->s0->a); // CHECK: 10 + printf ("%d\n", v->s0->b); // CHECK: 10 + printf ("%d\n", v->s0->c); // CHECK: 10 + printf ("%d\n", v->s1->a); // CHECK: 20 + printf ("%d\n", v->s1->b); // CHECK: 23 + printf ("%d\n", v->s1->c); // CHECK: 25 + printf ("%d\n", v->s2->a); // CHECK: 30 + printf ("%d\n", v->s2->b); // CHECK: 37 + printf ("%d\n", v->s2->c); // CHECK: 30 + + free(v->s0); + free(v->s1); + free(v->s2); + free(v); + + return 0; +} diff --git a/offload/test/mapping/chained_containing_structs_2.cc b/offload/test/mapping/chained_containing_structs_2.cc new file mode 100644 index 0000000000000..29c4c8b7fedfd --- /dev/null +++ b/offload/test/mapping/chained_containing_structs_2.cc @@ -0,0 +1,76 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic +// XFAIL: * + +#include +#include +#include + +struct R { + int d; + int e; + int f; +}; + +struct S { + R *r0; + R *r1; + R *r2; +}; + +struct T { + S *s0; + S *s1; + S *s2; +}; + +int main() { + T *v = (T *) malloc (sizeof(T)); + + v->s0 = (S *) malloc (sizeof(S)); + v->s1 = (S *) malloc (sizeof(S)); + v->s2 = (S *) malloc (sizeof(S)); + + v->s0->r0 = (R *) calloc (1, sizeof(R)); + v->s0->r1 = (R *) calloc (1, sizeof(R)); + v->s0->r2 = (R *) calloc (1, sizeof(R)); + + v->s1->r0 = (R *) calloc (1, sizeof(R)); + v->s1->r1 = (R *) calloc (1, sizeof(R)); + v->s1->r2 = (R *) calloc (1, sizeof(R)); + + v->s2->r0 = (R *) calloc (1, sizeof(R)); + v->s2->r1 = (R *) calloc (1, sizeof(R)); + v->s2->r2 = (R *) calloc (1, sizeof(R)); + + #pragma omp target map(to: v->s1, v->s2, *v->s1, v->s1->r1, *v->s2, v->s2->r0) \ + map(tofrom: v->s1->r1->d, v->s1->r1->e, v->s1->r2->d, v->s1->r2->f, v->s2->r0->e) + { + v->s1->r1->d += 3; + v->s1->r1->e += 5; + v->s1->r2->d += 7; + v->s1->r2->f += 9; + v->s2->r0->e += 11; + } + + printf ("%d\n", v->s1->r1->d); // CHECK: 3 + printf ("%d\n", v->s1->r1->e); // CHECK: 5 + printf ("%d\n", v->s1->r2->d); // CHECK: 7 + printf ("%d\n", v->s1->r2->f); // CHECK: 9 + printf ("%d\n", v->s2->r0->e); // CHECK: 11 + + free(v->s0->r0); + free(v->s0->r1); + free(v->s0->r2); + free(v->s1->r0); + free(v->s1->r1); + free(v->s1->r2); + free(v->s2->r0); + free(v->s2->r1); + free(v->s2->r2); + free(v->s0); + free(v->s1); + free(v->s2); + free(v); + + return 0; +} diff --git a/offload/test/mapping/chained_containing_structs_3.cc b/offload/test/mapping/chained_containing_structs_3.cc new file mode 100644 index 0000000000000..23555bf69110d --- /dev/null +++ b/offload/test/mapping/chained_containing_structs_3.cc @@ -0,0 +1,217 @@ +// RUN: %libomptarget-compilexx-run-and-check-generic + +// XFAIL: * + +#include +#include +#include +#include + +#include + +struct R { + int d; + int e; + int f; +}; + +struct S { + int a; + int b; + struct { + int c; + R r; + R *rp; + } sub; + int g; +}; + +struct T { + int a; + int *ptr; + int b; +}; + +int main() { + R r; + R *rp = new R; + S s; + S *sp = new S; + T t; + T *tp = new T; + + memset(&r, 0, sizeof(R)); + memset(rp, 0, sizeof(R)); + memset(&s, 0, sizeof(S)); + memset(sp, 0, sizeof(S)); + memset(&t, 0, sizeof(T)); + memset(tp, 0, sizeof(T)); + + s.sub.rp = new R; + sp->sub.rp = new R; + + memset(s.sub.rp, 0, sizeof(R)); + memset(sp->sub.rp, 0, sizeof(R)); + + t.ptr = new int[10]; + tp->ptr = new int[10]; + + memset(t.ptr, 0, sizeof(int)*10); + memset(tp->ptr, 0, sizeof(int)*10); + +#pragma omp target map(tofrom: r) map(tofrom: r.e) +{ + r.d++; + r.e += 2; + r.f += 3; +} + printf ("%d\n", r.d); // CHECK: 1 + printf ("%d\n", r.e); // CHECK-NEXT: 2 + printf ("%d\n", r.f); // CHECK-NEXT: 3 + +#pragma omp target map(tofrom: rp[:1]) map(tofrom: rp->e) +{ + rp->d++; + rp->e += 2; + rp->f += 3; +} + + printf ("%d\n", rp->d); // CHECK-NEXT: 1 + printf ("%d\n", rp->e); // CHECK-NEXT: 2 + printf ("%d\n", rp->f); // CHECK-NEXT: 3 + + int v; + int *orig_addr_v = &v; + bool separate_memory_space; + +#pragma omp target data map(v) + { + void *mapped_ptr_v = + omp_get_mapped_ptr(orig_addr_v, omp_get_default_device()); + separate_memory_space = mapped_ptr_v != (void*) orig_addr_v; + } + + const char *mapping_flavour = separate_memory_space ? "separate" : "unified"; + +#pragma omp target map(to: s) map(tofrom: s.sub.r.e) +{ + s.b++; + s.sub.r.d+=2; + s.sub.r.e+=3; + s.sub.r.f+=4; +} + + printf ("%d/%s\n", s.b, mapping_flavour); + printf ("%d/%s\n", s.sub.r.d, mapping_flavour); + printf ("%d/%s\n", s.sub.r.e, mapping_flavour); + printf ("%d/%s\n", s.sub.r.f, mapping_flavour); + + // CHECK: {{0/separate|1/unified}} + // CHECK-NEXT: {{0/separate|2/unified}} + // CHECK-NEXT: 3 + // CHECK-NEXT: {{0/separate|4/unified}} + +#pragma omp target map(to: s, s.b) map(to: s.sub.rp[:1]) map(tofrom: s.sub.rp->e) +{ + s.b++; + s.sub.rp->d+=2; + s.sub.rp->e+=3; + s.sub.rp->f+=4; +} + + printf ("%d/%s\n", s.b, mapping_flavour); + printf ("%d/%s\n", s.sub.rp->d, mapping_flavour); + printf ("%d/%s\n", s.sub.rp->e, mapping_flavour); + printf ("%d/%s\n", s.sub.rp->f, mapping_flavour); + + // CHECK-NEXT: {{0/separate|2/unified}} + // CHECK-NEXT: {{0/separate|2/unified}} + // CHECK-NEXT: 3 + // CHECK-NEXT: {{0/separate|4/unified}} + +#pragma omp target map(to: sp[:1]) map(tofrom: sp->sub.r.e) +{ + sp->b++; + sp->sub.r.d+=2; + sp->sub.r.e+=3; + sp->sub.r.f+=4; +} + + printf ("%d/%s\n", sp->b, mapping_flavour); + printf ("%d/%s\n", sp->sub.r.d, mapping_flavour); + printf ("%d/%s\n", sp->sub.r.e, mapping_flavour); + printf ("%d/%s\n", sp->sub.r.f, mapping_flavour); + + // CHECK-NEXT: {{0/separate|1/unified}} + // CHECK-NEXT: {{0/separate|2/unified}} + // CHECK-NEXT: 3 + // CHECK-NEXT: {{0/separate|4/unified}} + +#pragma omp target map(to: sp[:1]) map(to: sp->sub.rp[:1]) map(tofrom: sp->sub.rp->e) +{ + sp->b++; + sp->sub.rp->d+=2; + sp->sub.rp->e+=3; + sp->sub.rp->f+=4; +} + + printf ("%d/%s\n", sp->b, mapping_flavour); + printf ("%d/%s\n", sp->sub.rp->d, mapping_flavour); + printf ("%d/%s\n", sp->sub.rp->e, mapping_flavour); + printf ("%d/%s\n", sp->sub.rp->f, mapping_flavour); + + // CHECK-NEXT: {{0/separate|2/unified}} + // CHECK-NEXT: {{0/separate|2/unified}} + // CHECK-NEXT: 3 + // CHECK-NEXT: {{0/separate|4/unified}} + +#pragma omp target map(tofrom: t) map(tofrom: t.ptr[2:1]) +{ + t.a++; + t.ptr[2]+=2; + t.b+=3; +} + + printf ("%d\n", t.a); // CHECK-NEXT: 1 + printf ("%d\n", t.ptr[2]); // CHECK-NEXT: 2 + printf ("%d\n", t.b); // CHECK-NEXT: 3 + +#pragma omp target map(tofrom: t) map(tofrom: t.a) +{ + t.b++; +} + + printf ("%d\n", t.b); // CHECK-NEXT: 4 + +#pragma omp target map(tofrom: t) map(tofrom: t.ptr[2:1], t.a) +{ + t.a++; + t.ptr[2]+=2; + t.b+=3; +} + + printf ("%d\n", t.a); // CHECK-NEXT: 2 + printf ("%d\n", t.ptr[2]); // CHECK-NEXT: 4 + printf ("%d\n", t.b); // CHECK-NEXT: 7 + +#pragma omp target map(tofrom: t) map(tofrom: t.ptr[2:1], t.a) +{ + /* Empty */ +} + + printf ("%d\n", t.a); // CHECK-NEXT: 2 + printf ("%d\n", t.ptr[2]); // CHECK-NEXT: 4 + printf ("%d\n", t.b); // CHECK-NEXT: 7 + + delete s.sub.rp; + delete sp->sub.rp; + + delete[] t.ptr; + delete[] tp->ptr; + + delete rp; + delete sp; + delete tp; + + return 0; +}