[llvm] [OpenMP] Add tests for mapping of chained 'containing' structs (PR #156703)
Julian Brown via llvm-commits
llvm-commits at lists.llvm.org
Thu Sep 4 09:10:47 PDT 2025
https://github.com/jtb20 updated https://github.com/llvm/llvm-project/pull/156703
>From 9df11a3de47af2d43066c3a4f308f30762a439d2 Mon Sep 17 00:00:00 2001
From: Julian Brown <julian.brown at amd.com>
Date: Wed, 3 Sep 2025 05:34:34 -0500
Subject: [PATCH] [OpenMP] Add tests for mapping of chained 'containing'
structs
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
https://github.com/llvm/llvm-project/pull/153683.
---
.../mapping/chained_containing_structs_1.cc | 58 +++++
.../mapping/chained_containing_structs_2.cc | 76 ++++++
.../mapping/chained_containing_structs_3.cc | 217 ++++++++++++++++++
3 files changed, 351 insertions(+)
create mode 100644 offload/test/mapping/chained_containing_structs_1.cc
create mode 100644 offload/test/mapping/chained_containing_structs_2.cc
create mode 100644 offload/test/mapping/chained_containing_structs_3.cc
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 <cstdlib>
+#include <cstdio>
+#include <cassert>
+
+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 <cstdlib>
+#include <cstdio>
+#include <cassert>
+
+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 <cstdlib>
+#include <cstdio>
+#include <cassert>
+#include <cstring>
+
+#include <omp.h>
+
+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;
+}
More information about the llvm-commits
mailing list