[Mlir-commits] [clang] [llvm] [mlir] [WIP][OpenMP] Enable `ATTACH`-style maps for mappers. (PR #166874)
Abhinav Gaba
llvmlistbot at llvm.org
Thu May 14 19:02:18 PDT 2026
https://github.com/abhinavgaba updated https://github.com/llvm/llvm-project/pull/166874
>From 96dc2a4d9207dd807a9d792c7f265dc68b80dbf0 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Thu, 23 Apr 2026 14:07:35 -0700
Subject: [PATCH 01/13] This is a follow-up to #153683.
The change is experimental for now as it simply enables the
ATTACH-style maps, and updates a few libomptarget tests.
clang tests haven't been updated yet.
Also, we might need to add some new "implicit" maps. It's not clear
whether any additional implicit maps should be added for a
list-item mapped via a mapper, either default or non-default,
based on the contents of the mapper.
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 22 ++++++++++++++++++-
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 20 +++++++++++++++--
.../mapping/declare_mapper_nested_mappers.cpp | 7 +++---
.../test/mapping/declare_mapper_target.cpp | 2 +-
.../mapping/declare_mapper_target_data.cpp | 2 +-
.../declare_mapper_target_data_enter_exit.cpp | 2 +-
.../mapping/declare_mapper_target_update.cpp | 2 +-
7 files changed, 47 insertions(+), 10 deletions(-)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 332b439c87472..40d8080a34f9c 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -9383,7 +9383,27 @@ class MappableExprsHandler {
/// Constructor for the declare mapper directive.
MappableExprsHandler(const OMPDeclareMapperDecl &Dir, CodeGenFunction &CGF)
- : CurDir(&Dir), CGF(CGF), AttachPtrComparator(*this) {}
+ : CurDir(&Dir), CGF(CGF), AttachPtrComparator(*this) {
+ auto CollectAttachPtrExprsForClauseComponents = [this](const auto *C) {
+ for (auto L : C->component_lists()) {
+ OMPClauseMappableExprCommon::MappableExprComponentListRef Components =
+ std::get<1>(L);
+ if (!Components.empty())
+ collectAttachPtrExprInfo(Components, CurDir);
+ }
+ };
+
+ // Populate the AttachPtrExprMap for all component lists from map-related
+ // clauses in the declare mapper directive.
+ for (const auto *Cl : Dir.clauses()) {
+ if (const auto *C = dyn_cast<OMPMapClause>(Cl))
+ CollectAttachPtrExprsForClauseComponents(C);
+ else if (const auto *C = dyn_cast<OMPToClause>(Cl))
+ CollectAttachPtrExprsForClauseComponents(C);
+ else if (const auto *C = dyn_cast<OMPFromClause>(Cl))
+ CollectAttachPtrExprsForClauseComponents(C);
+ }
+ }
/// Generate code for the combined entry if we have a partially mapped struct
/// and take care of the mapping flags of the arguments corresponding to
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 125620bd49502..29bd501ed36e2 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -10047,8 +10047,24 @@ Expected<Function *> OpenMPIRBuilder::emitUserDefinedMapper(
CurMapType->addIncoming(FromMapType, FromBB);
CurMapType->addIncoming(MemberMapType, ToElseBB);
- Value *OffloadingArgs[] = {MapperHandle, CurBaseArg, CurBeginArg,
- CurSizeArg, CurMapType, CurNameArg};
+ // We need to propagete the DELETE bit to each map inserted by the mapper.
+ //
+ // OpenMP 6.0:281:34: The effect of the mapper modifier is to remove the
+ // list item from the map clause and to apply the clauses specified in the
+ // declared mapper to the construct on which the map clause appears...
+ // If any modifier with the map-type-modifying property appears in the map
+ // clause then the effect is as if that modifier appears in each map clause
+ // specified in the declared mapper.
+ Value *DeleteBitMask = Builder.CreateAnd(
+ MapType,
+ Builder.getInt64(
+ static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>(
+ OpenMPOffloadMappingFlags::OMP_MAP_DELETE)));
+ Value *CurMapTypeWithDelete =
+ Builder.CreateOr(CurMapType, DeleteBitMask, "omp.maptype.with.delete");
+
+ Value *OffloadingArgs[] = {MapperHandle, CurBaseArg, CurBeginArg,
+ CurSizeArg, CurMapTypeWithDelete, CurNameArg};
auto ChildMapperFn = CustomMapperCB(I);
if (!ChildMapperFn)
diff --git a/offload/test/mapping/declare_mapper_nested_mappers.cpp b/offload/test/mapping/declare_mapper_nested_mappers.cpp
index a59ed6980ec4c..5966bb6127af8 100644
--- a/offload/test/mapping/declare_mapper_nested_mappers.cpp
+++ b/offload/test/mapping/declare_mapper_nested_mappers.cpp
@@ -7,7 +7,8 @@ typedef struct {
int a;
double *b;
} C;
-#pragma omp declare mapper(id1 : C s) map(to : s.a) map(from : s.b[0 : 2])
+#pragma omp declare mapper(id1 : C s) map(to : s.a) map(alloc : s.b) \
+ map(from : s.b[0 : 2])
typedef struct {
int e;
@@ -16,7 +17,7 @@ typedef struct {
short *g;
} D;
#pragma omp declare mapper(default : D r) map(from : r.e) \
- map(mapper(id1), tofrom : r.f) map(tofrom : r.g[0 : r.h])
+ map(mapper(id1), tofrom : r.f) map(alloc : r.g) map(tofrom : r.g[0 : r.h])
int main() {
constexpr int N = 10;
@@ -56,7 +57,7 @@ int main() {
spp[0][0].f.b[1] = 40;
spp[0][0].g[1] = 50;
}
- printf("%d %d %d %d\n", spp00fa, spp00fb_r, spp00fg1, spp00fg_r);
+ printf("%d %d %d %d\n", spp00fa, spp00fb_r, spp00fg1, spp00fg_r);
// CHECK: 222 0 30 0
printf("%d %d %4.5f %d %d %d\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.b[1],
spp[0][0].f.b == &x[0] ? 1 : 0, spp[0][0].g[1],
diff --git a/offload/test/mapping/declare_mapper_target.cpp b/offload/test/mapping/declare_mapper_target.cpp
index 4d7237e94657a..c642949b2a367 100644
--- a/offload/test/mapping/declare_mapper_target.cpp
+++ b/offload/test/mapping/declare_mapper_target.cpp
@@ -10,7 +10,7 @@ class C {
int *a;
};
-#pragma omp declare mapper(id : C s) map(s.a[0 : NUM])
+#pragma omp declare mapper(id : C s) map(alloc : s.a) map(s.a[0 : NUM])
int main() {
C c;
diff --git a/offload/test/mapping/declare_mapper_target_data.cpp b/offload/test/mapping/declare_mapper_target_data.cpp
index 7f09844014005..fd2c8f4886401 100644
--- a/offload/test/mapping/declare_mapper_target_data.cpp
+++ b/offload/test/mapping/declare_mapper_target_data.cpp
@@ -10,7 +10,7 @@ class C {
int *a;
};
-#pragma omp declare mapper(id : C s) map(s.a[0 : NUM])
+#pragma omp declare mapper(id : C s) map(alloc : s.a) map(s.a[0 : NUM])
int main() {
C c;
diff --git a/offload/test/mapping/declare_mapper_target_data_enter_exit.cpp b/offload/test/mapping/declare_mapper_target_data_enter_exit.cpp
index f5fad8b8fe332..245462eab5785 100644
--- a/offload/test/mapping/declare_mapper_target_data_enter_exit.cpp
+++ b/offload/test/mapping/declare_mapper_target_data_enter_exit.cpp
@@ -10,7 +10,7 @@ class C {
int *a;
};
-#pragma omp declare mapper(id : C s) map(s.a[0 : NUM])
+#pragma omp declare mapper(id : C s) map(alloc : s.a) map(s.a[0 : NUM])
int main() {
C c;
diff --git a/offload/test/mapping/declare_mapper_target_update.cpp b/offload/test/mapping/declare_mapper_target_update.cpp
index fe4597b76908f..4e053e3e56139 100644
--- a/offload/test/mapping/declare_mapper_target_update.cpp
+++ b/offload/test/mapping/declare_mapper_target_update.cpp
@@ -10,7 +10,7 @@ class C {
int *a;
};
-#pragma omp declare mapper(id : C s) map(s.a[0 : NUM])
+#pragma omp declare mapper(id : C s) map(alloc : s.a) map(s.a[0 : NUM])
int main() {
C c;
>From 607393d36915da8937b89a6a9e1777c97cfa15ca Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Thu, 23 Apr 2026 15:26:28 -0700
Subject: [PATCH 02/13] [WIP] Add tests that require map-type bit propagation.
---
...apper_map_mbr_ptee_then_present_mbr_ptee.c | 44 +++++++++++++
.../mapper_map_mbr_then_present_mbr_ptee.c | 54 ++++++++++++++++
offload/test/mapping/mapper_map_ptee_only.c | 46 +++++++++++++
.../mapper_map_ptee_only_2_ptr_indirections.c | 58 +++++++++++++++++
.../mapping/mapper_map_ptee_only_2ndlevel.c | 50 +++++++++++++++
.../test/mapping/mapper_map_ptee_only_array.c | 64 +++++++++++++++++++
.../mapper_map_ptr_ptee_nomapper_del_ptee.c | 43 +++++++++++++
.../mapper_map_ptr_ptee_nomapper_del_ptr.c | 43 +++++++++++++
.../multiple_deletes_within_one_struct.c | 44 +++++++++++++
9 files changed, 446 insertions(+)
create mode 100644 offload/test/mapping/mapper_map_mbr_ptee_then_present_mbr_ptee.c
create mode 100644 offload/test/mapping/mapper_map_mbr_then_present_mbr_ptee.c
create mode 100644 offload/test/mapping/mapper_map_ptee_only.c
create mode 100644 offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections.c
create mode 100644 offload/test/mapping/mapper_map_ptee_only_2ndlevel.c
create mode 100644 offload/test/mapping/mapper_map_ptee_only_array.c
create mode 100644 offload/test/mapping/mapper_map_ptr_ptee_nomapper_del_ptee.c
create mode 100644 offload/test/mapping/mapper_map_ptr_ptee_nomapper_del_ptr.c
create mode 100644 offload/test/mapping/multiple_deletes_within_one_struct.c
diff --git a/offload/test/mapping/mapper_map_mbr_ptee_then_present_mbr_ptee.c b/offload/test/mapping/mapper_map_mbr_ptee_then_present_mbr_ptee.c
new file mode 100644
index 0000000000000..bed78e442c4ea
--- /dev/null
+++ b/offload/test/mapping/mapper_map_mbr_ptee_then_present_mbr_ptee.c
@@ -0,0 +1,44 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+// Check that it's ok to first map member of a struct and its pointee, and then
+// do a map(present) on a mapper that maps them internally.
+
+int x[10];
+
+typedef struct {
+ int x;
+ int dummy[10000];
+ int *p;
+} S;
+
+#pragma omp declare mapper(default : S s) map(s.x, s.p[0 : 10])
+
+S s1;
+
+void print_status(void *p, const char *name) {
+ int present = omp_target_is_present(p, omp_get_default_device());
+ printf("%s is %spresent\n", name, present ? "" : "not ");
+}
+
+int main() {
+ s1.p = (int *)&x;
+
+#pragma omp target enter data map(alloc : s1.x, s1.p[0:10])
+ printf("After mapping\n");
+ print_status(&s1.x, "x"); // CHECK: x is present
+ print_status(&s1.dummy, "dummy"); // CHECK: dummy is not present
+ print_status(&s1.p, "p"); // CHECK: p is not present
+ print_status(&s1.p[0], "p[0]"); // CHECK: p[0] is not present
+ printf("\n");
+
+ // This present check should pass.
+#pragma omp target enter data map(present, alloc : s1)
+ printf("After deleting\n");
+ print_status(&s1.x, "x"); // CHECK: x is not present
+ print_status(&s1.dummy, "dummy"); // CHECK: dummy is not present
+ print_status(&s1.p, "p"); // CHECK: p is not present
+ print_status(&s1.p[0], "p[0]"); // CHECK: p[0] is not present
+}
diff --git a/offload/test/mapping/mapper_map_mbr_then_present_mbr_ptee.c b/offload/test/mapping/mapper_map_mbr_then_present_mbr_ptee.c
new file mode 100644
index 0000000000000..9341d0572c4ba
--- /dev/null
+++ b/offload/test/mapping/mapper_map_mbr_then_present_mbr_ptee.c
@@ -0,0 +1,54 @@
+// RUN: %libomptarget-compile-generic
+// RUN: %libomptarget-run-fail-generic 2>&1 | %fcheck-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+// Check that the present check fails if we map a struct member, then do a
+// do a map(present) on a mapper that maps both a member and a pointee.
+
+int x[10];
+
+typedef struct {
+ int x;
+ int dummy[10000];
+ int *p;
+} S;
+
+#pragma omp declare mapper(default : S s) map(s.x, s.p[0 : 10])
+
+S s1;
+
+void print_status(void *p, const char *name) {
+ int present = omp_target_is_present(p, omp_get_default_device());
+ printf("%s is %spresent\n", name, present ? "" : "not ");
+}
+
+int main() {
+ s1.p = (int *)&x;
+
+ // CHECK: addr=0x[[#%x,HOST_ADDR:]], size=[[#%u,SIZE:]]
+ fprintf(stderr, "addr=%p, size=%ld\n", &s1.p[0], 10 * sizeof(s1.p[0]));
+
+#pragma omp target enter data map(alloc : s1.x)
+ printf("After mapping\n");
+ print_status(&s1.x, "x"); // CHECK: x is present
+ print_status(&s1.dummy, "dummy"); // CHECK: dummy is not present
+ print_status(&s1.p, "p"); // CHECK: p is not present
+ print_status(&s1.p[0], "p[0]"); // CHECK: p[0] is not present
+ printf("\n");
+
+ // This present check should fail!
+
+ // clang-format off
+ // CHECK: omptarget message: device mapping required by 'present' motion modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes)
+ // CHECK: omptarget fatal error 1: failure of target construct while offloading is mandatory
+ // clang-format on
+
+#pragma omp target enter data map(present, alloc : s1)
+ printf("After deleting\n");
+ print_status(&s1.x, "x");
+ print_status(&s1.dummy, "dummy");
+ print_status(&s1.p, "p");
+ print_status(&s1.p[0], "p[0]");
+}
diff --git a/offload/test/mapping/mapper_map_ptee_only.c b/offload/test/mapping/mapper_map_ptee_only.c
new file mode 100644
index 0000000000000..eec30ab6a372d
--- /dev/null
+++ b/offload/test/mapping/mapper_map_ptee_only.c
@@ -0,0 +1,46 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+// This test ensures that when mapping `s.x` and `s.p[0:10] via a mapper, the
+// storage of `s.p` and hence `s.dummy` is not unnecessarily allocated.
+
+int x[10];
+
+typedef struct {
+ int x;
+ int dummy[10000];
+ int *p;
+} S;
+
+#pragma omp declare mapper(default : S s) map(s.x, s.p[0 : 10])
+
+S s1;
+
+void print_status(void *p, const char *name) {
+ int present = omp_target_is_present(p, omp_get_default_device());
+ printf("%s is %spresent\n", name, present ? "" : "not ");
+}
+
+int main() {
+ s1.p = (int *)&x;
+
+#pragma omp target enter data map(alloc : s1)
+ printf("After mapping\n");
+ print_status(&s1.x, "x"); // CHECK: x is present
+ // FIXME: These will be resolved once we enable ATTACH style maps for mappers.
+ print_status(&s1.dummy, "dummy"); // EXPECTED: dummy is not present
+ // CHECK: dummy is present
+ print_status(&s1.p, "p"); // EXPECTED: p is not present
+ // CHECK: p is present
+ print_status(&s1.p[0], "p[0]"); // CHECK: p[0] is present
+ printf("\n");
+
+#pragma omp target exit data map(delete : s1)
+ printf("After deleting\n");
+ print_status(&s1.x, "x"); // CHECK: x is not present
+ print_status(&s1.dummy, "dummy"); // CHECK: dummy is not present
+ print_status(&s1.p, "p"); // CHECK: p is not present
+ print_status(&s1.p[0], "p[0]"); // CHECK: p[0] is not present
+}
diff --git a/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections.c b/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections.c
new file mode 100644
index 0000000000000..21500b1d2ad18
--- /dev/null
+++ b/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections.c
@@ -0,0 +1,58 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+// Check that the mapping of members of a struct and pointee data using a mapper
+// properly allocates/deletes all fields.
+
+int x[10];
+
+typedef struct {
+ int x;
+ int y;
+ int dummy[10000];
+ int *p;
+} S1;
+
+typedef struct {
+ S1 *s1p;
+ int z;
+} S2;
+
+#pragma omp declare mapper(default : S2 s2) \
+ map(s2.z, s2.s1p -> x, s2.s1p->y, s2.s1p->p[0 : 10])
+
+S1 s1;
+S2 s2;
+
+void print_status(void *p, const char *name) {
+ int present = omp_target_is_present(p, omp_get_default_device());
+ printf("%s is %spresent\n", name, present ? "" : "not ");
+}
+
+int main() {
+ s2.s1p = &s1;
+ s2.s1p->p = (int *)&x;
+
+#pragma omp target enter data map(alloc : s2)
+ printf("After mapping\n");
+ print_status(&s2.s1p->x, "x"); // CHECK: x is present
+ print_status(&s2.s1p->y, "y"); // CHECK: y is present
+ print_status(&s2.z, "z"); // CHECK: z is present
+ print_status(&s2.s1p->dummy, "dummy"); // CHECK: dummy is not present
+ // FIXME: These will be resolved once we enable ATTACH style maps for mappers.
+ print_status(&s2.s1p->p, "p"); // EXPECTED: p is not present
+ // CHECK: p is present
+ print_status(&s2.s1p->p[0], "p[0]"); // CHECK: p[0] is present
+ printf("\n");
+
+#pragma omp target exit data map(delete : s2)
+ printf("After deleting\n");
+ print_status(&s2.s1p->x, "x"); // CHECK: x is not present
+ print_status(&s2.s1p->y, "y"); // CHECK: y is not present
+ print_status(&s2.z, "z"); // CHECK: z is not present
+ print_status(&s2.s1p->dummy, "dummy"); // CHECK: dummy is not present
+ print_status(&s2.s1p->p, "p"); // CHECK: p is not present
+ print_status(&s2.s1p->p[0], "p[0]"); // CHECK: p[0] is not present
+}
diff --git a/offload/test/mapping/mapper_map_ptee_only_2ndlevel.c b/offload/test/mapping/mapper_map_ptee_only_2ndlevel.c
new file mode 100644
index 0000000000000..a522d6cc2cd44
--- /dev/null
+++ b/offload/test/mapping/mapper_map_ptee_only_2ndlevel.c
@@ -0,0 +1,50 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+// This test ensures that when mapping `s2.s1.x` and `s2.s1.p[0:10] via a
+// mapper, the storage of `s2.s1.p` and hence `s2.s1.dummy` is not unnecessarily
+// allocated.
+
+int x[10];
+
+typedef struct {
+ int x;
+ int dummy[10000];
+ int *p;
+} S1;
+
+typedef struct {
+ S1 s1;
+} S2;
+
+#pragma omp declare mapper(default : S2 s2) map(s2.s1.x, s2.s1.p[0 : 10])
+
+S2 s2;
+
+void print_status(void *p, const char *name) {
+ int present = omp_target_is_present(p, omp_get_default_device());
+ printf("%s is %spresent\n", name, present ? "" : "not ");
+}
+
+int main() {
+ s2.s1.p = (int *)&x;
+
+#pragma omp target enter data map(alloc : s2)
+ printf("After mapping\n");
+ print_status(&s2.s1.x, "x"); // CHECK: x is present
+ print_status(&s2.s1.dummy, "dummy"); // CHECK: dummy is not present
+ // FIXME: These will be resolved once we enable ATTACH style maps for mappers.
+ print_status(&s2.s1.p, "p"); // EXPECTED: p is not present
+ // CHECK: p is present
+ print_status(&s2.s1.p[0], "p[0]"); // CHECK: p[0] is present
+ printf("\n");
+
+#pragma omp target exit data map(delete : s2)
+ printf("After deleting\n");
+ print_status(&s2.s1.x, "x"); // CHECK: x is not present
+ print_status(&s2.s1.dummy, "dummy"); // CHECK: dummy is not present
+ print_status(&s2.s1.p, "p"); // CHECK: p is not present
+ print_status(&s2.s1.p[0], "p[0]"); // CHECK: p[0] is not present
+}
diff --git a/offload/test/mapping/mapper_map_ptee_only_array.c b/offload/test/mapping/mapper_map_ptee_only_array.c
new file mode 100644
index 0000000000000..65a5e32c2bcd3
--- /dev/null
+++ b/offload/test/mapping/mapper_map_ptee_only_array.c
@@ -0,0 +1,64 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+// Test that a mapper that maps a var and a pointee, works on an array of
+// structs in terms of properly allocating/deleting the data.
+
+int x[2][10];
+
+typedef struct {
+ int x;
+ int dummy[10000];
+ int *p;
+} S;
+
+#pragma omp declare mapper(default : S s) map(s.x, s.p[0 : 10])
+
+S s1[2];
+
+void print_status(void *p, const char *name) {
+ int present = omp_target_is_present(p, omp_get_default_device());
+ printf("%s is %spresent\n", name, present ? "" : "not ");
+}
+
+int main() {
+ s1[0].x = 111;
+ s1[1].x = 222;
+ s1[0].p = (int *)&x[0];
+ s1[1].p = (int *)&x[1];
+
+#pragma omp target enter data map(to : s1)
+ printf("After mapping\n");
+ print_status(&s1[0].x, "s1[0].x"); // CHECK: s1[0].x is present
+ // dummy/p being present is not ideal, but that's what we get with the
+ // current implementation because we need to map the full contiguous
+ // chunk for the array first before invoking the mapper.
+ print_status(&s1[0].dummy, "s1[0].dummy"); // CHECK: s1[0].dummy is present
+ print_status(&s1[0].p, "s1[0].p"); // CHECK: s1[0].p is present
+ print_status(&s1[0].p[0], "s1[0].p[0]"); // CHECK: s1[0].p[0] is present
+ print_status(&s1[1].x, "s1[1].x"); // CHECK: s1[1].x is present
+ print_status(&s1[1].dummy, "s1[1].dummy"); // CHECK: s1[1].dummy is present
+ print_status(&s1[1].p, "s1[1].p"); // CHECK: s1[1].p is present
+ print_status(&s1[1].p[0], "s1[1].p[0]"); // CHECK: s1[1].p[0] is present
+
+#pragma omp target map(present, alloc : s1)
+ {
+ printf("%d %d\n", s1[0].x, s1[1].x);
+ }
+
+ printf("\n");
+#pragma omp target exit data map(delete : s1)
+ printf("After deleting\n");
+ print_status(&s1[0].x, "s1[0].x"); // CHECK: s1[0].x is not present
+ print_status(&s1[0].dummy,
+ "s1[0].dummy"); // CHECK: s1[0].dummy is not present
+ print_status(&s1[0].p, "s1[0].p"); // CHECK: s1[0].p is not present
+ print_status(&s1[0].p[0], "s1[0].p[0]"); // CHECK: s1[0].p[0] is not present
+ print_status(&s1[1].x, "s1[1].x"); // CHECK: s1[1].x is not present
+ print_status(&s1[1].dummy,
+ "s1[1].dummy"); // CHECK: s1[1].dummy is not present
+ print_status(&s1[1].p, "s1[1].p"); // CHECK: s1[1].p is not present
+ print_status(&s1[1].p[0], "s1[1].p[0]"); // CHECK: s1[1].p[0] is not present
+}
diff --git a/offload/test/mapping/mapper_map_ptr_ptee_nomapper_del_ptee.c b/offload/test/mapping/mapper_map_ptr_ptee_nomapper_del_ptee.c
new file mode 100644
index 0000000000000..26a7098016207
--- /dev/null
+++ b/offload/test/mapping/mapper_map_ptr_ptee_nomapper_del_ptee.c
@@ -0,0 +1,43 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+// This test ensures that when mapping `s.p` and `s.p[0:10] via a mapper, they
+// occupy different storage blocks in memory, so deleting `s.p[0:10]` does not
+// delete `s.p`.
+
+int x[10];
+
+typedef struct {
+ int *p;
+} S;
+
+#pragma omp declare mapper(default : S s) map(s.p, s.p[0 : 10])
+
+S s1;
+
+void print_status(void *p, const char *name) {
+ int present = omp_target_is_present(p, omp_get_default_device());
+ printf("%s is %spresent\n", name, present ? "" : "not ");
+}
+int main() {
+ s1.p = (int *)&x;
+
+#pragma omp target enter data map(alloc : s1)
+ printf("After mapping ptr and ptee\n");
+ print_status(&s1.p, "p"); // CHECK: p is present
+ print_status(&s1.p[0], "p[0]"); // CHECK: p[0] is present
+ printf("\n");
+
+#pragma omp target exit data map(delete : s1.p[0 : 10])
+ printf("After deleting ptee\n");
+ print_status(&s1.p, "p"); // CHECK: p is present
+ print_status(&s1.p[0], "p[0]"); // CHECK: p[0] is not present
+ printf("\n");
+
+#pragma omp target exit data map(delete : s1.p)
+ printf("After deleting ptr\n");
+ print_status(&s1.p, "p"); // CHECK: p is not present
+ print_status(&s1.p[0], "p[0]"); // CHECK: p[0] is not present
+}
diff --git a/offload/test/mapping/mapper_map_ptr_ptee_nomapper_del_ptr.c b/offload/test/mapping/mapper_map_ptr_ptee_nomapper_del_ptr.c
new file mode 100644
index 0000000000000..7457a988e5933
--- /dev/null
+++ b/offload/test/mapping/mapper_map_ptr_ptee_nomapper_del_ptr.c
@@ -0,0 +1,43 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+// This test ensures that when mapping `s.p` and `s.p[0:10] via a mapper, they
+// occupy different storage blocks in memory, so deleting `s.p` does not delete
+// `s.p[0:10]`.
+
+int x[10];
+
+typedef struct {
+ int *p;
+} S;
+
+#pragma omp declare mapper(default : S s) map(s.p, s.p[0 : 10])
+
+S s1;
+
+void print_status(void *p, const char *name) {
+ int present = omp_target_is_present(p, omp_get_default_device());
+ printf("%s is %spresent\n", name, present ? "" : "not ");
+}
+int main() {
+ s1.p = (int *)&x;
+
+#pragma omp target enter data map(alloc : s1)
+ printf("After mapping ptr and ptee\n");
+ print_status(&s1.p, "p"); // CHECK: p is present
+ print_status(&s1.p[0], "p[0]"); // CHECK: p[0] is present
+ printf("\n");
+
+#pragma omp target exit data map(delete : s1.p)
+ printf("After deleting ptr\n");
+ print_status(&s1.p, "p"); // CHECK: p is not present
+ print_status(&s1.p[0], "p[0]"); // CHECK: p[0] is present
+ printf("\n");
+
+#pragma omp target exit data map(delete : s1.p[0 : 10])
+ printf("After deleting ptee\n");
+ print_status(&s1.p, "p"); // CHECK: p is not present
+ print_status(&s1.p[0], "p[0]"); // CHECK: p[0] is not present
+}
diff --git a/offload/test/mapping/multiple_deletes_within_one_struct.c b/offload/test/mapping/multiple_deletes_within_one_struct.c
new file mode 100644
index 0000000000000..d5292a2acc0f7
--- /dev/null
+++ b/offload/test/mapping/multiple_deletes_within_one_struct.c
@@ -0,0 +1,44 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+// This test ensures that when mapping `s.x` and `s.p[0:10], the
+// storage of `s.p` and `s.dummy` is not unnecessarily allocated.
+
+// Secondly, after deleting s1.x and s1.p[0], they are no longer
+// present on the device.
+
+int g[10];
+
+typedef struct {
+ int x;
+ int dummy[10000];
+ int *p;
+} S;
+
+S s1;
+
+void print_status(void *p, const char *name) {
+ int present = omp_target_is_present(p, omp_get_default_device());
+ printf("%s is %spresent\n", name, present ? "" : "not ");
+}
+
+int main() {
+ s1.p = (int *)&g;
+
+#pragma omp target enter data map(alloc : s1.x, s1.p[0 : 10])
+ printf("After mapping\n");
+ print_status(&s1.x, "x"); // CHECK: x is present
+ print_status(&s1.dummy, "dummy"); // CHECK: dummy is not present
+ print_status(&s1.p, "p"); // CHECK: p is not present
+ print_status(&s1.p[0], "p[0]"); // CHECK: p[0] is present
+ printf("\n");
+
+#pragma omp target exit data map(delete : s1.x) map(delete : s1.p[0 : 10])
+ printf("After deleting\n");
+ print_status(&s1.x, "x"); // CHECK: x is not present
+ print_status(&s1.dummy, "dummy"); // CHECK: dummy is not present
+ print_status(&s1.p, "p"); // CHECK: p is not present
+ print_status(&s1.p[0], "p[0]"); // CHECK: p[0] is not present
+}
>From a1a6afc23f7790530085b809e7bd2e2a3ca3bdea Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Thu, 23 Apr 2026 21:42:21 -0700
Subject: [PATCH 03/13] Add more tests.
---
offload/test/mapping/mapper_map_ptee_only.c | 9 +--
.../mapper_map_ptee_only_2_ptr_indirections.c | 6 +-
...r_map_ptee_only_2_ptr_indirections_array.c | 72 ++++++++++++++++++
.../mapping/mapper_map_ptee_only_2ndlevel.c | 6 +-
.../mapper_map_ptee_only_2ndlevel_array.c | 76 +++++++++++++++++++
.../test/mapping/mapper_map_ptee_only_array.c | 5 --
6 files changed, 155 insertions(+), 19 deletions(-)
create mode 100644 offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections_array.c
create mode 100644 offload/test/mapping/mapper_map_ptee_only_2ndlevel_array.c
diff --git a/offload/test/mapping/mapper_map_ptee_only.c b/offload/test/mapping/mapper_map_ptee_only.c
index eec30ab6a372d..3944545b6e23c 100644
--- a/offload/test/mapping/mapper_map_ptee_only.c
+++ b/offload/test/mapping/mapper_map_ptee_only.c
@@ -28,12 +28,9 @@ int main() {
#pragma omp target enter data map(alloc : s1)
printf("After mapping\n");
- print_status(&s1.x, "x"); // CHECK: x is present
- // FIXME: These will be resolved once we enable ATTACH style maps for mappers.
- print_status(&s1.dummy, "dummy"); // EXPECTED: dummy is not present
- // CHECK: dummy is present
- print_status(&s1.p, "p"); // EXPECTED: p is not present
- // CHECK: p is present
+ print_status(&s1.x, "x"); // CHECK: x is present
+ print_status(&s1.dummy, "dummy"); // CHECK: dummy is not present
+ print_status(&s1.p, "p"); // CHECK: p is not present
print_status(&s1.p[0], "p[0]"); // CHECK: p[0] is present
printf("\n");
diff --git a/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections.c b/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections.c
index 21500b1d2ad18..112f60165361f 100644
--- a/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections.c
+++ b/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections.c
@@ -41,10 +41,8 @@ int main() {
print_status(&s2.s1p->y, "y"); // CHECK: y is present
print_status(&s2.z, "z"); // CHECK: z is present
print_status(&s2.s1p->dummy, "dummy"); // CHECK: dummy is not present
- // FIXME: These will be resolved once we enable ATTACH style maps for mappers.
- print_status(&s2.s1p->p, "p"); // EXPECTED: p is not present
- // CHECK: p is present
- print_status(&s2.s1p->p[0], "p[0]"); // CHECK: p[0] is present
+ print_status(&s2.s1p->p, "p"); // CHECK: p is not present
+ print_status(&s2.s1p->p[0], "p[0]"); // CHECK: p[0] is present
printf("\n");
#pragma omp target exit data map(delete : s2)
diff --git a/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections_array.c b/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections_array.c
new file mode 100644
index 0000000000000..621b78ba17061
--- /dev/null
+++ b/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections_array.c
@@ -0,0 +1,72 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+// Array variant of mapper_map_ptee_only_2_ptr_indirections.c.
+// Exercises the nested-pointer-chain case (s2.s1p->p[0:10]) in an array
+// context, where inner MEMBER_OF bits must be shifted correctly and outer
+// MEMBER_OF must not be applied to the pointee entries.
+
+int x[2][10];
+
+typedef struct {
+ int x;
+ int y;
+ int dummy[10000];
+ int *p;
+} S1;
+
+typedef struct {
+ S1 *s1p;
+ int z;
+} S2;
+
+#pragma omp declare mapper(default : S2 s2) \
+ map(s2.z, s2.s1p->x, s2.s1p->y, s2.s1p->p[0 : 10])
+
+S1 s1arr[2];
+S2 s2arr[2];
+
+void print_status(void *p, const char *name) {
+ int present = omp_target_is_present(p, omp_get_default_device());
+ printf("%s is %spresent\n", name, present ? "" : "not ");
+}
+
+int main() {
+ s2arr[0].s1p = &s1arr[0];
+ s2arr[1].s1p = &s1arr[1];
+ s2arr[0].s1p->p = (int *)&x[0];
+ s2arr[1].s1p->p = (int *)&x[1];
+
+#pragma omp target enter data map(alloc : s2arr)
+ printf("After mapping\n");
+ print_status(&s2arr[0].s1p->x, "s2arr[0].x"); // CHECK: s2arr[0].x is present
+ print_status(&s2arr[0].s1p->y, "s2arr[0].y"); // CHECK: s2arr[0].y is present
+ print_status(&s2arr[0].z, "s2arr[0].z"); // CHECK: s2arr[0].z is present
+ print_status(&s2arr[0].s1p->dummy, "s2arr[0].dummy"); // CHECK: s2arr[0].dummy is not present
+ print_status(&s2arr[0].s1p->p, "s2arr[0].p"); // CHECK: s2arr[0].p is not present
+ print_status(&s2arr[0].s1p->p[0], "s2arr[0].p[0]"); // CHECK: s2arr[0].p[0] is present
+ print_status(&s2arr[1].s1p->x, "s2arr[1].x"); // CHECK: s2arr[1].x is present
+ print_status(&s2arr[1].s1p->y, "s2arr[1].y"); // CHECK: s2arr[1].y is present
+ print_status(&s2arr[1].z, "s2arr[1].z"); // CHECK: s2arr[1].z is present
+ print_status(&s2arr[1].s1p->dummy, "s2arr[1].dummy"); // CHECK: s2arr[1].dummy is not present
+ print_status(&s2arr[1].s1p->p, "s2arr[1].p"); // CHECK: s2arr[1].p is not present
+ print_status(&s2arr[1].s1p->p[0], "s2arr[1].p[0]"); // CHECK: s2arr[1].p[0] is present
+ printf("\n");
+
+#pragma omp target exit data map(delete : s2arr)
+ printf("After deleting\n");
+ print_status(&s2arr[0].s1p->x, "s2arr[0].x"); // CHECK: s2arr[0].x is not present
+ print_status(&s2arr[0].s1p->y, "s2arr[0].y"); // CHECK: s2arr[0].y is not present
+ print_status(&s2arr[0].z, "s2arr[0].z"); // CHECK: s2arr[0].z is not present
+ print_status(&s2arr[0].s1p->dummy, "s2arr[0].dummy"); // CHECK: s2arr[0].dummy is not present
+ print_status(&s2arr[0].s1p->p, "s2arr[0].p"); // CHECK: s2arr[0].p is not present
+ print_status(&s2arr[0].s1p->p[0], "s2arr[0].p[0]"); // CHECK: s2arr[0].p[0] is not present
+ print_status(&s2arr[1].s1p->x, "s2arr[1].x"); // CHECK: s2arr[1].x is not present
+ print_status(&s2arr[1].s1p->y, "s2arr[1].y"); // CHECK: s2arr[1].y is not present
+ print_status(&s2arr[1].z, "s2arr[1].z"); // CHECK: s2arr[1].z is not present
+ print_status(&s2arr[1].s1p->dummy, "s2arr[1].dummy"); // CHECK: s2arr[1].dummy is not present
+ print_status(&s2arr[1].s1p->p, "s2arr[1].p"); // CHECK: s2arr[1].p is not present
+ print_status(&s2arr[1].s1p->p[0], "s2arr[1].p[0]"); // CHECK: s2arr[1].p[0] is not present
+}
diff --git a/offload/test/mapping/mapper_map_ptee_only_2ndlevel.c b/offload/test/mapping/mapper_map_ptee_only_2ndlevel.c
index a522d6cc2cd44..2d5404a6d0cea 100644
--- a/offload/test/mapping/mapper_map_ptee_only_2ndlevel.c
+++ b/offload/test/mapping/mapper_map_ptee_only_2ndlevel.c
@@ -35,10 +35,8 @@ int main() {
printf("After mapping\n");
print_status(&s2.s1.x, "x"); // CHECK: x is present
print_status(&s2.s1.dummy, "dummy"); // CHECK: dummy is not present
- // FIXME: These will be resolved once we enable ATTACH style maps for mappers.
- print_status(&s2.s1.p, "p"); // EXPECTED: p is not present
- // CHECK: p is present
- print_status(&s2.s1.p[0], "p[0]"); // CHECK: p[0] is present
+ print_status(&s2.s1.p, "p"); // CHECK: p is not present
+ print_status(&s2.s1.p[0], "p[0]"); // CHECK: p[0] is present
printf("\n");
#pragma omp target exit data map(delete : s2)
diff --git a/offload/test/mapping/mapper_map_ptee_only_2ndlevel_array.c b/offload/test/mapping/mapper_map_ptee_only_2ndlevel_array.c
new file mode 100644
index 0000000000000..6a7ee78c4544a
--- /dev/null
+++ b/offload/test/mapping/mapper_map_ptee_only_2ndlevel_array.c
@@ -0,0 +1,76 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+// Test that a mapper on a nested struct maps the right members when applied to
+// an array of structs: s.x and s.p[0:10] are mapped; s.dummy and s.p itself
+// are not (modulo attach FIXME).
+
+int x[2][10];
+
+typedef struct {
+ int x;
+ int dummy[10000];
+ int *p;
+} S1;
+
+typedef struct {
+ S1 s1;
+} S2;
+
+#pragma omp declare mapper(default : S2 s2) map(s2.s1.x, s2.s1.p[0 : 10])
+
+S2 s2arr[2];
+
+void print_status(void *p, const char *name) {
+ int present = omp_target_is_present(p, omp_get_default_device());
+ printf("%s is %spresent\n", name, present ? "" : "not ");
+}
+
+int main() {
+ s2arr[0].s1.x = 111;
+ s2arr[1].s1.x = 222;
+ s2arr[0].s1.p = (int *)&x[0];
+ s2arr[1].s1.p = (int *)&x[1];
+
+#pragma omp target enter data map(to : s2arr)
+ printf("After mapping\n");
+ print_status(&s2arr[0].s1.x, "s2arr[0].s1.x"); // CHECK: s2arr[0].s1.x is present
+ // dummy/p being present is not ideal, but that's what we get with the
+ // current implementation because we need to map the full contiguous
+ // chunk for the array first before invoking the mapper.
+ print_status(&s2arr[0].s1.dummy,
+ "s2arr[0].s1.dummy"); // CHECK: s2arr[0].s1.dummy is present
+ print_status(&s2arr[0].s1.p,
+ "s2arr[0].s1.p"); // CHECK: s2arr[0].s1.p is present
+ print_status(&s2arr[0].s1.p[0],
+ "s2arr[0].s1.p[0]"); // CHECK: s2arr[0].s1.p[0] is present
+ print_status(&s2arr[1].s1.x, "s2arr[1].s1.x"); // CHECK: s2arr[1].s1.x is present
+ print_status(&s2arr[1].s1.dummy,
+ "s2arr[1].s1.dummy"); // CHECK: s2arr[1].s1.dummy is present
+ print_status(&s2arr[1].s1.p,
+ "s2arr[1].s1.p"); // CHECK: s2arr[1].s1.p is present
+ print_status(&s2arr[1].s1.p[0],
+ "s2arr[1].s1.p[0]"); // CHECK: s2arr[1].s1.p[0] is present
+
+ printf("\n");
+#pragma omp target exit data map(delete : s2arr)
+ printf("After deleting\n");
+ print_status(&s2arr[0].s1.x,
+ "s2arr[0].s1.x"); // CHECK: s2arr[0].s1.x is not present
+ print_status(&s2arr[0].s1.dummy,
+ "s2arr[0].s1.dummy"); // CHECK: s2arr[0].s1.dummy is not present
+ print_status(&s2arr[0].s1.p,
+ "s2arr[0].s1.p"); // CHECK: s2arr[0].s1.p is not present
+ print_status(&s2arr[0].s1.p[0],
+ "s2arr[0].s1.p[0]"); // CHECK: s2arr[0].s1.p[0] is not present
+ print_status(&s2arr[1].s1.x,
+ "s2arr[1].s1.x"); // CHECK: s2arr[1].s1.x is not present
+ print_status(&s2arr[1].s1.dummy,
+ "s2arr[1].s1.dummy"); // CHECK: s2arr[1].s1.dummy is not present
+ print_status(&s2arr[1].s1.p,
+ "s2arr[1].s1.p"); // CHECK: s2arr[1].s1.p is not present
+ print_status(&s2arr[1].s1.p[0],
+ "s2arr[1].s1.p[0]"); // CHECK: s2arr[1].s1.p[0] is not present
+}
diff --git a/offload/test/mapping/mapper_map_ptee_only_array.c b/offload/test/mapping/mapper_map_ptee_only_array.c
index 65a5e32c2bcd3..c5f194d4b6c21 100644
--- a/offload/test/mapping/mapper_map_ptee_only_array.c
+++ b/offload/test/mapping/mapper_map_ptee_only_array.c
@@ -43,11 +43,6 @@ int main() {
print_status(&s1[1].p, "s1[1].p"); // CHECK: s1[1].p is present
print_status(&s1[1].p[0], "s1[1].p[0]"); // CHECK: s1[1].p[0] is present
-#pragma omp target map(present, alloc : s1)
- {
- printf("%d %d\n", s1[0].x, s1[1].x);
- }
-
printf("\n");
#pragma omp target exit data map(delete : s1)
printf("After deleting\n");
>From 68643ae0e07aa844ce4a7400ed1f74df4afc3c36 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Thu, 23 Apr 2026 21:46:58 -0700
Subject: [PATCH 04/13] Use attach-style maps for mappers and fix mapper
map-type propagation.
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 67 +++++++++++++-
.../llvm/Frontend/OpenMP/OMPIRBuilder.h | 7 ++
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 92 ++++++++++++++++++-
.../Frontend/OpenMPIRBuilderTest.cpp | 1 +
.../OpenMP/OpenMPToLLVMIRTranslation.cpp | 5 +
5 files changed, 162 insertions(+), 10 deletions(-)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 332b439c87472..767deefdb76d7 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7583,6 +7583,7 @@ class MappableExprsHandler {
AttachInfo.AttachPteeAddr.emitRawPointer(CGF));
CombinedInfo.Sizes.push_back(PointerSize);
CombinedInfo.Types.push_back(OpenMPOffloadMappingFlags::OMP_MAP_ATTACH);
+ CombinedInfo.DontAddMemberOfInMapper.push_back(false);
CombinedInfo.Mappers.push_back(nullptr);
CombinedInfo.NonContigInfo.Dims.push_back(1);
}
@@ -7684,6 +7685,7 @@ class MappableExprsHandler {
CombinedInfo.Sizes.push_back(
CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/false));
CombinedInfo.Types.push_back(Flags);
+ CombinedInfo.DontAddMemberOfInMapper.push_back(false);
CombinedInfo.Mappers.push_back(nullptr);
CombinedInfo.NonContigInfo.Dims.push_back(IsNonContiguous ? DimSize : 1);
}
@@ -8372,10 +8374,14 @@ class MappableExprsHandler {
}
}
- if (!IsMappingWholeStruct)
+ if (!IsMappingWholeStruct) {
CombinedInfo.Types.push_back(Flags);
- else
+ CombinedInfo.DontAddMemberOfInMapper.push_back(HasAttachPtr);
+ } else {
StructBaseCombinedInfo.Types.push_back(Flags);
+ StructBaseCombinedInfo.DontAddMemberOfInMapper.push_back(
+ HasAttachPtr);
+ }
}
// If we have encountered a member expression so far, keep track of the
@@ -8976,6 +8982,7 @@ class MappableExprsHandler {
if (HasUdpFbNullify)
Flags |= OpenMPOffloadMappingFlags::OMP_MAP_FB_NULLIFY;
UseDeviceDataCombinedInfo.Types.push_back(Flags);
+ UseDeviceDataCombinedInfo.DontAddMemberOfInMapper.push_back(false);
UseDeviceDataCombinedInfo.Mappers.push_back(nullptr);
};
@@ -9383,7 +9390,27 @@ class MappableExprsHandler {
/// Constructor for the declare mapper directive.
MappableExprsHandler(const OMPDeclareMapperDecl &Dir, CodeGenFunction &CGF)
- : CurDir(&Dir), CGF(CGF), AttachPtrComparator(*this) {}
+ : CurDir(&Dir), CGF(CGF), AttachPtrComparator(*this) {
+ auto CollectAttachPtrExprsForClauseComponents = [this](const auto *C) {
+ for (auto L : C->component_lists()) {
+ OMPClauseMappableExprCommon::MappableExprComponentListRef Components =
+ std::get<1>(L);
+ if (!Components.empty())
+ collectAttachPtrExprInfo(Components, CurDir);
+ }
+ };
+
+ // Populate the AttachPtrExprMap for all component lists from map-related
+ // clauses in the declare mapper directive.
+ for (const auto *Cl : Dir.clauses()) {
+ if (const auto *C = dyn_cast<OMPMapClause>(Cl))
+ CollectAttachPtrExprsForClauseComponents(C);
+ else if (const auto *C = dyn_cast<OMPToClause>(Cl))
+ CollectAttachPtrExprsForClauseComponents(C);
+ else if (const auto *C = dyn_cast<OMPFromClause>(Cl))
+ CollectAttachPtrExprsForClauseComponents(C);
+ }
+ }
/// Generate code for the combined entry if we have a partially mapped struct
/// and take care of the mapping flags of the arguments corresponding to
@@ -9457,6 +9484,14 @@ class MappableExprsHandler {
: !PartialStruct.PreliminaryMapData.BasePointers.empty()
? OpenMPOffloadMappingFlags::OMP_MAP_PTR_AND_OBJ
: OpenMPOffloadMappingFlags::OMP_MAP_TARGET_PARAM);
+ // Combined entries with an attach pointer occupy different storage than
+ // the original item, so if this map is for a mapper, it should not be a
+ // member of the original item being mapped. e.g.:
+ // map(s2.s1p->x, s2.s1p->y)
+ // combined entry:
+ // s2.s1p[0], s2.s1p->x, sizeof(s1p->x..y), ALLOC
+ // This occupies different storage than s2.
+ CombinedInfo.DontAddMemberOfInMapper.push_back(AttachInfo.isValid());
// If any element has the present modifier, then make sure the runtime
// doesn't attempt to allocate the struct.
if (CurTypes.end() !=
@@ -9572,6 +9607,7 @@ class MappableExprsHandler {
OpenMPOffloadMappingFlags::OMP_MAP_LITERAL |
OpenMPOffloadMappingFlags::OMP_MAP_MEMBER_OF |
OpenMPOffloadMappingFlags::OMP_MAP_IMPLICIT);
+ CombinedInfo.DontAddMemberOfInMapper.push_back(false);
CombinedInfo.Mappers.push_back(nullptr);
}
for (const LambdaCapture &LC : RD->captures()) {
@@ -9612,6 +9648,7 @@ class MappableExprsHandler {
OpenMPOffloadMappingFlags::OMP_MAP_LITERAL |
OpenMPOffloadMappingFlags::OMP_MAP_MEMBER_OF |
OpenMPOffloadMappingFlags::OMP_MAP_IMPLICIT);
+ CombinedInfo.DontAddMemberOfInMapper.push_back(false);
CombinedInfo.Mappers.push_back(nullptr);
}
}
@@ -9904,6 +9941,7 @@ class MappableExprsHandler {
CurCaptureVarInfo.Types.push_back(
OpenMPOffloadMappingFlags::OMP_MAP_LITERAL |
OpenMPOffloadMappingFlags::OMP_MAP_TARGET_PARAM);
+ CurCaptureVarInfo.DontAddMemberOfInMapper.push_back(false);
CurCaptureVarInfo.Mappers.push_back(nullptr);
return;
}
@@ -10302,6 +10340,7 @@ class MappableExprsHandler {
if (IsImplicit)
CombinedInfo.Types.back() |= OpenMPOffloadMappingFlags::OMP_MAP_IMPLICIT;
+ CombinedInfo.DontAddMemberOfInMapper.push_back(false);
// No user-defined mapper for default mapping.
CombinedInfo.Mappers.push_back(nullptr);
}
@@ -10518,14 +10557,30 @@ getNestedDistributeDirective(ASTContext &Ctx, const OMPExecutableDirective &D) {
/// size*sizeof(Ty), clearToFromMember(type));
/// // Map members.
/// for (unsigned i = 0; i < size; i++) {
+/// N = __tgt_mapper_num_components(rt_mapper_handle);
/// // For each component specified by this mapper:
/// for (auto c : begin[i]->all_components) {
+/// // MEMBER_OF grouping: tie this component to the current array element
+/// // (component N) by adding N<<48. Exceptions:
+/// // - ATTACH entries are not members of any struct storage range.
+/// // - Pointee entries (reached via a pointer member) occupy separate
+/// // storage; their inner MEMBER_OF bits are shifted by N instead.
+/// if (c.isAttach() || c.isPointee())
+/// member_type = c.arg_type + (c.hasInnerMemberOf() ? N<<48 : 0);
+/// else
+/// member_type = c.arg_type + N<<48;
+/// // Map-type-modifying bits (ALWAYS, DELETE, CLOSE, PRESENT) from the
+/// // outer map clause are propagated to each component, except ATTACH
+/// // entries (ATTACH|ALWAYS is reserved for attach(always), and other
+/// // modifier bits have no meaning for ATTACH).
+/// effective_type = c.isAttach() ? member_type
+/// : member_type | modifierBits(type);
/// if (c.hasMapper())
/// (*c.Mapper())(rt_mapper_handle, c.arg_base, c.arg_begin, c.arg_size,
-/// c.arg_type, c.arg_name);
+/// effective_type, c.arg_name);
/// else
/// __tgt_push_mapper_component(rt_mapper_handle, c.arg_base,
-/// c.arg_begin, c.arg_size, c.arg_type,
+/// c.arg_begin, c.arg_size, effective_type,
/// c.arg_name);
/// }
/// }
@@ -10741,6 +10796,7 @@ static void genMapInfoForCaptures(
CurInfo.Types.push_back(OpenMPOffloadMappingFlags::OMP_MAP_LITERAL |
OpenMPOffloadMappingFlags::OMP_MAP_TARGET_PARAM |
OpenMPOffloadMappingFlags::OMP_MAP_IMPLICIT);
+ CurInfo.DontAddMemberOfInMapper.push_back(false);
CurInfo.Mappers.push_back(nullptr);
} else {
const ValueDecl *CapturedVD =
@@ -10898,6 +10954,7 @@ static void emitTargetCallKernelLaunch(
CombinedInfo.Sizes.push_back(CGF.Builder.getInt64(0));
CombinedInfo.Types.push_back(OpenMPOffloadMappingFlags::OMP_MAP_TARGET_PARAM |
OpenMPOffloadMappingFlags::OMP_MAP_LITERAL);
+ CombinedInfo.DontAddMemberOfInMapper.push_back(false);
if (!CombinedInfo.Names.empty())
CombinedInfo.Names.push_back(NullPtr);
CombinedInfo.Exprs.push_back(nullptr);
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
index dbd8f0c6b8927..1120b80512b84 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
@@ -2821,6 +2821,7 @@ class OpenMPIRBuilder {
using MapNamesArrayTy = SmallVector<Constant *, 4>;
using MapDimArrayTy = SmallVector<uint64_t, 4>;
using MapNonContiguousArrayTy = SmallVector<MapValuesArrayTy, 4>;
+ using MapSkipMemberOfArrayTy = SmallVector<bool, 4>;
/// This structure contains combined information generated for mappable
/// clauses, including base pointers, pointers, sizes, map types, user-defined
@@ -2839,6 +2840,10 @@ class OpenMPIRBuilder {
MapValuesArrayTy Sizes;
MapFlagsArrayTy Types;
MapNamesArrayTy Names;
+ /// When true, emitUserDefinedMapper must not add a new
+ /// outer MEMBER_OF for this entry (unless there's already a MEMBER_OF set,
+ /// in which case it gets shifted/adjusted).
+ MapSkipMemberOfArrayTy DontAddMemberOfInMapper;
StructNonContiguousInfo NonContigInfo;
/// Append arrays in \a CurInfo.
@@ -2851,6 +2856,8 @@ class OpenMPIRBuilder {
Sizes.append(CurInfo.Sizes.begin(), CurInfo.Sizes.end());
Types.append(CurInfo.Types.begin(), CurInfo.Types.end());
Names.append(CurInfo.Names.begin(), CurInfo.Names.end());
+ DontAddMemberOfInMapper.append(CurInfo.DontAddMemberOfInMapper.begin(),
+ CurInfo.DontAddMemberOfInMapper.end());
NonContigInfo.Dims.append(CurInfo.NonContigInfo.Dims.begin(),
CurInfo.NonContigInfo.Dims.end());
NonContigInfo.Offsets.append(CurInfo.NonContigInfo.Offsets.begin(),
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 125620bd49502..e5b731b784c0a 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -9964,12 +9964,67 @@ Expected<Function *> OpenMPIRBuilder::emitUserDefinedMapper(
? Info->Names[I]
: Constant::getNullValue(Builder.getPtrTy());
- // Extract the MEMBER_OF field from the map type.
Value *OriMapType = Builder.getInt64(
static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>(
Info->Types[I]));
- Value *MemberMapType =
- Builder.CreateNUWAdd(OriMapType, ShiftedPreviousSize);
+ auto RawType =
+ static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>(
+ Info->Types[I]);
+ constexpr uint64_t MemberOfMask = 0xffff000000000000ULL;
+ bool IsAttach =
+ RawType &
+ static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>(
+ OpenMPOffloadMappingFlags::OMP_MAP_ATTACH);
+
+ // Add MEMBER_OF (ShiftedPreviousSize) to link this sub-map with the
+ // current array element. Two cases must not get a new outer MEMBER_OF:
+ //
+ // 1. Entries marked as DontAddMemberOfInMapper. These represent entries
+ // that represent the pointee data that does not occupy the same
+ // storage space as the array-element being mapped.
+ //
+ // 2. ATTACH entries: These do not represent an actual map, but pointer
+ // attachment, which is deferred to the end of the region on which the
+ // map occurs. It is not involved in any allocation/deallocation or
+ // ref-count tracking.
+ //
+ // Example 1:
+ // mapper: #pragma omp declare mapper(id: S s) map(s.x, s.p[0:10])
+ // use: S arr[2]; ... map(arr)
+ // entries per element:
+ //
+ // &arr[i], &arr[i].x, sizeof(int), MEMBER_OF(N)|TO|FROM
+ // &arr[i].p[0], &arr[i].p[0], 10*sizeof(int), TO|FROM (*)
+ // &arr[i].p, &arr[i].p[0], sizeof(int*), ATTACH
+ // (*) DontAddMemberOfInMapper=true: s.p[0:10] occupies different
+ // storage than s.
+ //
+ // Example 2: Struct with multiple pointees that share pointer-member storage
+ // mapper: #pragma omp declare mapper(S2 s2) map(s2.z, s2.s1p->x, s2.s1p->y)
+ // use: S2 arr[2]; ... map(arr)
+ // entries per element:
+ //
+ // &arr[i], &arr[i].z, sizeof(int), MEMBER_OF(N)|TO|FROM
+ // &arr[i].s1p[0], &arr[i].s1p->x, sizeof(s1p->x..y), ALLOC (**)
+ // &arr[i].s1p[0], &arr[i].s1p->x, sizeof(int), MEMBER_OF(N+2)|TO|FROM
+ // &arr[i].s1p[0], &arr[i].s1p->y, sizeof(int), MEMBER_OF(N+2)|TO|FROM
+ // &arr[i].s1p, &arr[i].s1p->x, sizeof(ptr), ATTACH
+ // (**) DontAddMemberOfInMapper=true (set by emitCombinedEntry): the *s1p
+ // pointee occupies separate storage from arr[i]. The individual x/y
+ // entries carry inner MEMBER_OF(2) bits which are shifted by N instead
+ // of getting a new outer layer.
+ // (N = __tgt_mapper_num_components() at loop body start)
+ Value *MemberMapType;
+ if (IsAttach) {
+ MemberMapType = OriMapType;
+ } else if (Info->DontAddMemberOfInMapper[I]) {
+ if (RawType & MemberOfMask)
+ MemberMapType = Builder.CreateNUWAdd(OriMapType, ShiftedPreviousSize);
+ else
+ MemberMapType = OriMapType;
+ } else {
+ MemberMapType = Builder.CreateNUWAdd(OriMapType, ShiftedPreviousSize);
+ }
// Combine the map type inherited from user-defined mapper with that
// specified in the program. According to the OMP_MAP_TO and OMP_MAP_FROM
@@ -10047,8 +10102,35 @@ Expected<Function *> OpenMPIRBuilder::emitUserDefinedMapper(
CurMapType->addIncoming(FromMapType, FromBB);
CurMapType->addIncoming(MemberMapType, ToElseBB);
- Value *OffloadingArgs[] = {MapperHandle, CurBaseArg, CurBeginArg,
- CurSizeArg, CurMapType, CurNameArg};
+ // Propagate map-type-modifying bits from the outer map clause to each map
+ // inserted by the mapper.
+ //
+ // OpenMP 6.0:281:34: The effect of the mapper modifier is to remove the
+ // list item from the map clause and to apply the clauses specified in the
+ // declared mapper to the construct on which the map clause appears...
+ // If any modifier with the map-type-modifying property appears in the map
+ // clause then the effect is as if that modifier appears in each map clause
+ // specified in the declared mapper.
+ //
+ // Map-type-modifying bits: ALWAYS, DELETE, CLOSE, PRESENT.
+ Value *ModifierBitMask = Builder.CreateAnd(
+ MapType,
+ Builder.getInt64(
+ static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>(
+ OpenMPOffloadMappingFlags::OMP_MAP_ALWAYS |
+ OpenMPOffloadMappingFlags::OMP_MAP_DELETE |
+ OpenMPOffloadMappingFlags::OMP_MAP_CLOSE |
+ OpenMPOffloadMappingFlags::OMP_MAP_PRESENT)));
+ Value *CurMapTypeWithModifiers = Builder.CreateOr(
+ CurMapType, ModifierBitMask, "omp.maptype.with.modifiers");
+
+ // ATTACH entries must not receive map-type-modifying bits: ATTACH|ALWAYS is
+ // reserved for the attach(always) map-type modifier, and other modifier bits
+ // (DELETE, CLOSE, PRESENT) have no meaning for an ATTACH entry.
+ Value *FinalMapType = IsAttach ? CurMapType : CurMapTypeWithModifiers;
+
+ Value *OffloadingArgs[] = {MapperHandle, CurBaseArg, CurBeginArg,
+ CurSizeArg, FinalMapType, CurNameArg};
auto ChildMapperFn = CustomMapperCB(I);
if (!ChildMapperFn)
diff --git a/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp b/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp
index 5b9b280ec7671..e2b7f4fb99de0 100644
--- a/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp
+++ b/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp
@@ -8266,6 +8266,7 @@ TEST_F(OpenMPIRBuilderTest, EmitOffloadingArraysNonContigCountExpression) {
CombinedInfo.Types.push_back(static_cast<omp::OpenMPOffloadMappingFlags>(
omp::OpenMPOffloadMappingFlags::OMP_MAP_NON_CONTIG |
omp::OpenMPOffloadMappingFlags::OMP_MAP_TO));
+ CombinedInfo.DontAddMemberOfInMapper.push_back(false);
CombinedInfo.Names.push_back(
Builder.CreateGlobalString("data", "data_name", 0, M.get()));
diff --git a/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp b/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp
index 8614aed1ab80c..5bb97d8fd6588 100644
--- a/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp
+++ b/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp
@@ -5160,6 +5160,7 @@ static void collectMapDataFromMapOperands(
mapData.BaseType.back(), builder, moduleTranslation));
mapData.MapClause.push_back(mapOp.getOperation());
mapData.Types.push_back(convertClauseMapFlags(mapOp.getMapType()));
+ mapData.DontAddMemberOfInMapper.push_back(false);
mapData.Names.push_back(LLVM::createMappingInformation(
mapOp.getLoc(), *moduleTranslation.getOpenMPBuilder()));
mapData.DevicePointers.push_back(llvm::OpenMPIRBuilder::DeviceInfoTy::None);
@@ -5210,6 +5211,7 @@ static void collectMapDataFromMapOperands(
mapData.MapClause.push_back(mapOp.getOperation());
mapData.Types.push_back(
llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_RETURN_PARAM);
+ mapData.DontAddMemberOfInMapper.push_back(false);
mapData.Names.push_back(LLVM::createMappingInformation(
mapOp.getLoc(), *moduleTranslation.getOpenMPBuilder()));
mapData.DevicePointers.push_back(devInfoTy);
@@ -5248,6 +5250,7 @@ static void collectMapDataFromMapOperands(
// rematerialized, so the address of the decriptor for a given object
// may change from one place to another.
mapData.Types.push_back(mapType);
+ mapData.DontAddMemberOfInMapper.push_back(false);
// Technically it's possible for a non-descriptor mapping to have
// both has-device-addr and ALWAYS, so lookup the mapper in case it
// exists.
@@ -5264,6 +5267,7 @@ static void collectMapDataFromMapOperands(
mapData.Types.push_back(
isDevicePtr ? mapType
: llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_LITERAL);
+ mapData.DontAddMemberOfInMapper.push_back(false);
mapData.Mappers.push_back(nullptr);
}
mapData.Names.push_back(LLVM::createMappingInformation(
@@ -7279,6 +7283,7 @@ convertOmpTarget(Operation &opInst, llvm::IRBuilderBase &builder,
combinedInfos.Types.push_back(
llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_TARGET_PARAM |
llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_LITERAL);
+ combinedInfos.DontAddMemberOfInMapper.push_back(false);
if (!combinedInfos.Names.empty())
combinedInfos.Names.push_back(nullPtr);
combinedInfos.Mappers.push_back(nullptr);
>From 64794270f721839f4e071427a8df2d60f1ba9362 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Fri, 24 Apr 2026 01:26:12 -0700
Subject: [PATCH 05/13] fix mbr_ptee present modifier checks.
---
.../mapping/mapper_map_mbr_ptee_then_present_mbr_ptee.c | 4 +++-
.../test/mapping/mapper_map_mbr_then_present_mbr_ptee.c | 9 ++-------
2 files changed, 5 insertions(+), 8 deletions(-)
diff --git a/offload/test/mapping/mapper_map_mbr_ptee_then_present_mbr_ptee.c b/offload/test/mapping/mapper_map_mbr_ptee_then_present_mbr_ptee.c
index bed78e442c4ea..00a8ba4bb4a81 100644
--- a/offload/test/mapping/mapper_map_mbr_ptee_then_present_mbr_ptee.c
+++ b/offload/test/mapping/mapper_map_mbr_ptee_then_present_mbr_ptee.c
@@ -31,11 +31,13 @@ int main() {
print_status(&s1.x, "x"); // CHECK: x is present
print_status(&s1.dummy, "dummy"); // CHECK: dummy is not present
print_status(&s1.p, "p"); // CHECK: p is not present
- print_status(&s1.p[0], "p[0]"); // CHECK: p[0] is not present
+ print_status(&s1.p[0], "p[0]"); // CHECK: p[0] is present
printf("\n");
// This present check should pass.
#pragma omp target enter data map(present, alloc : s1)
+
+#pragma omp target exit data map(delete: s1)
printf("After deleting\n");
print_status(&s1.x, "x"); // CHECK: x is not present
print_status(&s1.dummy, "dummy"); // CHECK: dummy is not present
diff --git a/offload/test/mapping/mapper_map_mbr_then_present_mbr_ptee.c b/offload/test/mapping/mapper_map_mbr_then_present_mbr_ptee.c
index 9341d0572c4ba..0939aa0b2330f 100644
--- a/offload/test/mapping/mapper_map_mbr_then_present_mbr_ptee.c
+++ b/offload/test/mapping/mapper_map_mbr_then_present_mbr_ptee.c
@@ -21,7 +21,7 @@ S s1;
void print_status(void *p, const char *name) {
int present = omp_target_is_present(p, omp_get_default_device());
- printf("%s is %spresent\n", name, present ? "" : "not ");
+ fprintf(stderr, "%s is %spresent\n", name, present ? "" : "not ");
}
int main() {
@@ -41,14 +41,9 @@ int main() {
// This present check should fail!
// clang-format off
- // CHECK: omptarget message: device mapping required by 'present' motion modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes)
+ // CHECK: omptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes)
// CHECK: omptarget fatal error 1: failure of target construct while offloading is mandatory
// clang-format on
#pragma omp target enter data map(present, alloc : s1)
- printf("After deleting\n");
- print_status(&s1.x, "x");
- print_status(&s1.dummy, "dummy");
- print_status(&s1.p, "p");
- print_status(&s1.p[0], "p[0]");
}
>From b12586f26c0bda593b397709c923c8497bdfce59 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Fri, 24 Apr 2026 02:07:50 -0700
Subject: [PATCH 06/13] Rename variable
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 15 +++--
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 67 +++++++++++------------
2 files changed, 42 insertions(+), 40 deletions(-)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 767deefdb76d7..f341e0b40801c 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7583,7 +7583,9 @@ class MappableExprsHandler {
AttachInfo.AttachPteeAddr.emitRawPointer(CGF));
CombinedInfo.Sizes.push_back(PointerSize);
CombinedInfo.Types.push_back(OpenMPOffloadMappingFlags::OMP_MAP_ATTACH);
- CombinedInfo.DontAddMemberOfInMapper.push_back(false);
+ // ATTACH entries don't participate in ref-count tracking, and must note be
+ // grouped as a member of the enclosing struct in emitUserDefinedMapper.
+ CombinedInfo.DontAddMemberOfInMapper.push_back(true);
CombinedInfo.Mappers.push_back(nullptr);
CombinedInfo.NonContigInfo.Dims.push_back(1);
}
@@ -8376,6 +8378,9 @@ class MappableExprsHandler {
if (!IsMappingWholeStruct) {
CombinedInfo.Types.push_back(Flags);
+ // Pointee entries (HasAttachPtr=true) occupy different storage than
+ // the pointer variable; emitUserDefinedMapper must not set
+ // MEMBER_OF for them.
CombinedInfo.DontAddMemberOfInMapper.push_back(HasAttachPtr);
} else {
StructBaseCombinedInfo.Types.push_back(Flags);
@@ -9484,13 +9489,13 @@ class MappableExprsHandler {
: !PartialStruct.PreliminaryMapData.BasePointers.empty()
? OpenMPOffloadMappingFlags::OMP_MAP_PTR_AND_OBJ
: OpenMPOffloadMappingFlags::OMP_MAP_TARGET_PARAM);
- // Combined entries with an attach pointer occupy different storage than
- // the original item, so if this map is for a mapper, it should not be a
- // member of the original item being mapped. e.g.:
+ // Combined entries with an attach pointer also occupy a different storage
+ // space than the pointer itself, so emitUserDefinedMapper should not add
+ // MEMBER_OF flags linking them. e.g.:
// map(s2.s1p->x, s2.s1p->y)
// combined entry:
// s2.s1p[0], s2.s1p->x, sizeof(s1p->x..y), ALLOC
- // This occupies different storage than s2.
+ // s2.s1p[0] occupies different storage than s2.s1p or s2.
CombinedInfo.DontAddMemberOfInMapper.push_back(AttachInfo.isValid());
// If any element has the present modifier, then make sure the runtime
// doesn't attempt to allocate the struct.
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index e5b731b784c0a..993456b3db28c 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -9971,22 +9971,11 @@ Expected<Function *> OpenMPIRBuilder::emitUserDefinedMapper(
static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>(
Info->Types[I]);
constexpr uint64_t MemberOfMask = 0xffff000000000000ULL;
- bool IsAttach =
- RawType &
- static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>(
- OpenMPOffloadMappingFlags::OMP_MAP_ATTACH);
- // Add MEMBER_OF (ShiftedPreviousSize) to link this sub-map with the
- // current array element. Two cases must not get a new outer MEMBER_OF:
- //
- // 1. Entries marked as DontAddMemberOfInMapper. These represent entries
- // that represent the pointee data that does not occupy the same
- // storage space as the array-element being mapped.
- //
- // 2. ATTACH entries: These do not represent an actual map, but pointer
- // attachment, which is deferred to the end of the region on which the
- // map occurs. It is not involved in any allocation/deallocation or
- // ref-count tracking.
+ // Add MEMBER_OF (ShiftedPreviousSize) to group this sub-map with the
+ // current array element (N = __tgt_mapper_num_components() at loop body
+ // start).
+
//
// Example 1:
// mapper: #pragma omp declare mapper(id: S s) map(s.x, s.p[0:10])
@@ -9994,30 +9983,34 @@ Expected<Function *> OpenMPIRBuilder::emitUserDefinedMapper(
// entries per element:
//
// &arr[i], &arr[i].x, sizeof(int), MEMBER_OF(N)|TO|FROM
- // &arr[i].p[0], &arr[i].p[0], 10*sizeof(int), TO|FROM (*)
- // &arr[i].p, &arr[i].p[0], sizeof(int*), ATTACH
- // (*) DontAddMemberOfInMapper=true: s.p[0:10] occupies different
- // storage than s.
+ // &arr[i].p[0], &arr[i].p[0], 10*sizeof(int), TO|FROM (*)
+ // &arr[i].p, &arr[i].p[0], sizeof(int*), ATTACH (**)
//
- // Example 2: Struct with multiple pointees that share pointer-member storage
- // mapper: #pragma omp declare mapper(S2 s2) map(s2.z, s2.s1p->x, s2.s1p->y)
+ // Example 2:
+ // mapper: #pragma omp declare mapper(S2 s2) map(s2.z, s2.s1p->x,
+ // s2.s1p->y)
// use: S2 arr[2]; ... map(arr)
// entries per element:
//
- // &arr[i], &arr[i].z, sizeof(int), MEMBER_OF(N)|TO|FROM
- // &arr[i].s1p[0], &arr[i].s1p->x, sizeof(s1p->x..y), ALLOC (**)
- // &arr[i].s1p[0], &arr[i].s1p->x, sizeof(int), MEMBER_OF(N+2)|TO|FROM
- // &arr[i].s1p[0], &arr[i].s1p->y, sizeof(int), MEMBER_OF(N+2)|TO|FROM
- // &arr[i].s1p, &arr[i].s1p->x, sizeof(ptr), ATTACH
- // (**) DontAddMemberOfInMapper=true (set by emitCombinedEntry): the *s1p
- // pointee occupies separate storage from arr[i]. The individual x/y
- // entries carry inner MEMBER_OF(2) bits which are shifted by N instead
- // of getting a new outer layer.
- // (N = __tgt_mapper_num_components() at loop body start)
+ // &arr[i], &arr[i].z, sizeof(int), MEMBER_OF(N)|TO|FROM
+ // &arr[i].s1p[0], &arr[i].s1p->x, sizeof(s1p->x..y), ALLOC (*)
+ // &arr[i].s1p[0], &arr[i].s1p->x, sizeof(int), MEMBER_OF(N+2)|TO|FROM
+ // &arr[i].s1p[0], &arr[i].s1p->y, sizeof(int), MEMBER_OF(N+2)|TO|FROM
+ // &arr[i].s1p, &arr[i].s1p->x, sizeof(ptr), ATTACH (**)
+ //
+ // x/y carry inner MEMBER_OF(2)
+ // which is shifted by N to become MEMBER_OF(N+2).
+ //
+ // Entries with DontAddMemberOfInMapper=true must not receive a
+ // new outer MEMBER_OF. They occupy a different storage block than the
+ // the enclosing struct (like pointee data (*)), or are ATTACH entries that
+ // represent pointer-attachment (**), and don't contribute to any ref-count
+ // entries for pointer members).
+ //
+ // If such an entry already has its own MEMBER_OF bits (like for s1p->x/y
+ // above), they are still shifted by N.
Value *MemberMapType;
- if (IsAttach) {
- MemberMapType = OriMapType;
- } else if (Info->DontAddMemberOfInMapper[I]) {
+ if (Info->DontAddMemberOfInMapper[I]) {
if (RawType & MemberOfMask)
MemberMapType = Builder.CreateNUWAdd(OriMapType, ShiftedPreviousSize);
else
@@ -10127,7 +10120,11 @@ Expected<Function *> OpenMPIRBuilder::emitUserDefinedMapper(
// ATTACH entries must not receive map-type-modifying bits: ATTACH|ALWAYS is
// reserved for the attach(always) map-type modifier, and other modifier bits
// (DELETE, CLOSE, PRESENT) have no meaning for an ATTACH entry.
- Value *FinalMapType = IsAttach ? CurMapType : CurMapTypeWithModifiers;
+ constexpr uint64_t AttachBit =
+ static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>(
+ OpenMPOffloadMappingFlags::OMP_MAP_ATTACH);
+ Value *FinalMapType =
+ (RawType & AttachBit) ? CurMapType : CurMapTypeWithModifiers;
Value *OffloadingArgs[] = {MapperHandle, CurBaseArg, CurBeginArg,
CurSizeArg, FinalMapType, CurNameArg};
>From 1b2803ca746492588d1e913f438ed59837c62d3f Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Fri, 24 Apr 2026 02:25:52 -0700
Subject: [PATCH 07/13] Add codegen test
---
...t_map_nested_ptr_member_mapper_codegen.cpp | 223 ++++++++++++++++++
1 file changed, 223 insertions(+)
create mode 100644 clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp
diff --git a/clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp b/clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp
new file mode 100644
index 0000000000000..d90d4fa3fcfbb
--- /dev/null
+++ b/clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp
@@ -0,0 +1,223 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --functions ".*mapper.*" --function-signature --check-globals --filter-out-after "getelem.*kernel" --filter-out "= alloca.*" --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _ --global-value-regex "\.offload_.*" --global-hex-value-regex ".offload_maptypes.*"
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
+
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+// &arr[i], &arr[i].z, sizeof(int), MEMBER_OF(N) | TO | FROM
+// &arr[i].s1p[0], &arr[i].s1p->x, sizeof(s1p->x..y), ALLOC (*)
+// &arr[i].s1p[0], &arr[i].s1p->x, sizeof(int), MEMBER_OF(N+2) | TO | FROM (**)
+// &arr[i].s1p[0], &arr[i].s1p->y, sizeof(int), MEMBER_OF(N+2) | TO | FROM (**)
+// &arr[i].s1p, &arr[i].s1p->x, sizeof(ptr), ATTACH (***)
+// (*) Pointee (combined entry): No MEMBER_OF addition
+// (**) Inner MEMBER_OF(2) entries: Shifted by N.
+// (***) ATTACH entry: No MEMBER_OF bit; no inherited map-type-modifying bits
+// (DELETE/ALWAYS/CLOSE/PRESENT).
+// (i = array element index; N = __tgt_mapper_num_components()):
+
+typedef struct {
+ int x;
+ int y;
+} S1;
+
+typedef struct {
+ S1 *s1p;
+ int z;
+} S2;
+
+#pragma omp declare mapper(default : S2 s2) map(s2.z, s2.s1p->x, s2.s1p->y)
+
+void foo(S2 *arr) {
+ // &arr, &arr[0], 2*sizeof(S2), TARGET_PARAM | TO
+ // (mapper handles individual members)
+#pragma omp target enter data map(to: arr[0:2])
+ {}
+}
+
+#endif
+//.
+// CHECK: @.offload_sizes = private unnamed_addr constant [1 x i64] [i64 32]
+// CHECK: @.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 [[#0x1]]]
+//.
+// CHECK-LABEL: define {{[^@]+}}@_Z3fooP2S2
+// CHECK-SAME: (ptr noundef [[ARR:%.*]]) #[[ATTR0:[0-9]+]] {
+// CHECK: entry:
+// CHECK: store ptr [[ARR]], ptr [[ARR_ADDR:%.*]], align 8
+// CHECK: [[TMP0:%.*]] = load ptr, ptr [[ARR_ADDR]], align 8
+// CHECK: [[TMP1:%.*]] = load ptr, ptr [[ARR_ADDR]], align 8
+// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds [[STRUCT_S2:%.*]], ptr [[TMP1]], i64 0
+// CHECK: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0
+// CHECK: store ptr [[TMP0]], ptr [[TMP2]], align 8
+// CHECK: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0
+// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP3]], align 8
+// CHECK: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0
+// CHECK: store ptr @.omp_mapper._ZTS2S2.default, ptr [[TMP4]], align 8
+// CHECK: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// CHECK: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// CHECK: call void @__tgt_target_data_begin_mapper(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 1, ptr [[TMP5]], ptr [[TMP6]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr [[DOTOFFLOAD_MAPPERS]])
+// CHECK: ret void
+//
+//
+// CHECK-LABEL: define {{[^@]+}}@.omp_mapper._ZTS2S2.default
+// CHECK-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]], ptr noundef [[TMP2:%.*]], i64 noundef [[TMP3:%.*]], i64 noundef [[TMP4:%.*]], ptr noundef [[TMP5:%.*]]) #[[ATTR1:[0-9]+]] {
+// CHECK: entry:
+// CHECK: store ptr [[TMP0]], ptr [[DOTADDR:%.*]], align 8
+// CHECK: store ptr [[TMP1]], ptr [[DOTADDR1:%.*]], align 8
+// CHECK: store ptr [[TMP2]], ptr [[DOTADDR2:%.*]], align 8
+// CHECK: store i64 [[TMP3]], ptr [[DOTADDR3:%.*]], align 8
+// CHECK: store i64 [[TMP4]], ptr [[DOTADDR4:%.*]], align 8
+// CHECK: store ptr [[TMP5]], ptr [[DOTADDR5:%.*]], align 8
+// CHECK: [[TMP6:%.*]] = load i64, ptr [[DOTADDR3]], align 8
+// CHECK: [[TMP7:%.*]] = load ptr, ptr [[DOTADDR]], align 8
+// CHECK: [[TMP8:%.*]] = load ptr, ptr [[DOTADDR1]], align 8
+// CHECK: [[TMP9:%.*]] = load ptr, ptr [[DOTADDR2]], align 8
+// CHECK: [[TMP10:%.*]] = udiv exact i64 [[TMP6]], 16
+// CHECK: [[TMP11:%.*]] = getelementptr [[STRUCT_S2:%.*]], ptr [[TMP9]], i64 [[TMP10]]
+// CHECK: [[TMP12:%.*]] = load i64, ptr [[DOTADDR4]], align 8
+// CHECK: [[TMP13:%.*]] = load ptr, ptr [[DOTADDR5]], align 8
+// CHECK: [[OMP_ARRAYINIT_ISARRAY:%.*]] = icmp sgt i64 [[TMP10]], 1
+// CHECK: [[TMP14:%.*]] = and i64 [[TMP12]], 8
+// CHECK: [[TMP15:%.*]] = icmp ne ptr [[TMP8]], [[TMP9]]
+// CHECK: [[TMP16:%.*]] = and i64 [[TMP12]], 16
+// CHECK: [[TMP17:%.*]] = icmp ne i64 [[TMP16]], 0
+// CHECK: [[TMP18:%.*]] = and i1 [[TMP15]], [[TMP17]]
+// CHECK: [[TMP19:%.*]] = or i1 [[OMP_ARRAYINIT_ISARRAY]], [[TMP18]]
+// CHECK: [[DOTOMP_ARRAY__INIT__DELETE:%.*]] = icmp eq i64 [[TMP14]], 0
+// CHECK: [[TMP20:%.*]] = and i1 [[TMP19]], [[DOTOMP_ARRAY__INIT__DELETE]]
+// CHECK: br i1 [[TMP20]], label [[DOTOMP_ARRAY__INIT:%.*]], label [[OMP_ARRAYMAP_HEAD:%.*]]
+// CHECK: .omp.array..init:
+// CHECK: [[TMP21:%.*]] = mul nuw i64 [[TMP10]], 16
+// CHECK: [[TMP22:%.*]] = and i64 [[TMP12]], -4
+// CHECK: [[TMP23:%.*]] = or i64 [[TMP22]], 512
+// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP7]], ptr [[TMP8]], ptr [[TMP9]], i64 [[TMP21]], i64 [[TMP23]], ptr [[TMP13]])
+// CHECK: br label [[OMP_ARRAYMAP_HEAD]]
+// CHECK: omp.arraymap.head:
+// CHECK: [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP9]], [[TMP11]]
+// CHECK: br i1 [[OMP_ARRAYMAP_ISEMPTY]], label [[OMP_DONE:%.*]], label [[OMP_ARRAYMAP_BODY:%.*]]
+// CHECK: omp.arraymap.body:
+// CHECK: [[OMP_ARRAYMAP_PTRCURRENT:%.*]] = phi ptr [ [[TMP9]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END28:%.*]] ]
+// CHECK: [[Z:%.*]] = getelementptr inbounds [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 1
+// CHECK: [[S1P:%.*]] = getelementptr inbounds [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
+// CHECK: [[S1P6:%.*]] = getelementptr inbounds [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
+// CHECK: [[TMP24:%.*]] = load ptr, ptr [[S1P6]], align 8
+// CHECK: [[X:%.*]] = getelementptr inbounds [[STRUCT_S1:%.*]], ptr [[TMP24]], i32 0, i32 0
+// CHECK: [[S1P7:%.*]] = getelementptr inbounds [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
+// CHECK: [[S1P8:%.*]] = getelementptr inbounds [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
+// CHECK: [[TMP25:%.*]] = load ptr, ptr [[S1P8]], align 8
+// CHECK: [[Y:%.*]] = getelementptr inbounds [[STRUCT_S1]], ptr [[TMP25]], i32 0, i32 1
+// CHECK: [[TMP26:%.*]] = getelementptr i32, ptr [[Z]], i32 1
+// CHECK: [[TMP27:%.*]] = ptrtoint ptr [[TMP26]] to i64
+// CHECK: [[TMP28:%.*]] = ptrtoint ptr [[S1P]] to i64
+// CHECK: [[TMP29:%.*]] = sub i64 [[TMP27]], [[TMP28]]
+// CHECK: [[TMP30:%.*]] = sdiv exact i64 [[TMP29]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)
+// CHECK: [[TMP31:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP7]])
+// CHECK: [[TMP32:%.*]] = shl i64 [[TMP31]], 48
+// CHECK: [[TMP33:%.*]] = add nuw i64 0, [[TMP32]]
+// CHECK: [[TMP34:%.*]] = and i64 [[TMP12]], 3
+// CHECK: [[TMP35:%.*]] = icmp eq i64 [[TMP34]], 0
+// CHECK: br i1 [[TMP35]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]]
+// CHECK: omp.type.alloc:
+// CHECK: [[TMP36:%.*]] = and i64 [[TMP33]], -4
+// CHECK: br label [[OMP_TYPE_END:%.*]]
+// CHECK: omp.type.alloc.else:
+// CHECK: [[TMP37:%.*]] = icmp eq i64 [[TMP34]], 1
+// CHECK: br i1 [[TMP37]], label [[OMP_TYPE_TO:%.*]], label [[OMP_TYPE_TO_ELSE:%.*]]
+// CHECK: omp.type.to:
+// CHECK: [[TMP38:%.*]] = and i64 [[TMP33]], -3
+// CHECK: br label [[OMP_TYPE_END]]
+// CHECK: omp.type.to.else:
+// CHECK: [[TMP39:%.*]] = icmp eq i64 [[TMP34]], 2
+// CHECK: br i1 [[TMP39]], label [[OMP_TYPE_FROM:%.*]], label [[OMP_TYPE_END]]
+// CHECK: omp.type.from:
+// CHECK: [[TMP40:%.*]] = and i64 [[TMP33]], -2
+// CHECK: br label [[OMP_TYPE_END]]
+// CHECK: omp.type.end:
+// CHECK: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP36]], [[OMP_TYPE_ALLOC]] ], [ [[TMP38]], [[OMP_TYPE_TO]] ], [ [[TMP40]], [[OMP_TYPE_FROM]] ], [ [[TMP33]], [[OMP_TYPE_TO_ELSE]] ]
+// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP7]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[S1P]], i64 [[TMP30]], i64 [[OMP_MAPTYPE]], ptr null)
+// CHECK: [[TMP41:%.*]] = add nuw i64 281474976710659, [[TMP32]]
+// CHECK: [[TMP42:%.*]] = and i64 [[TMP12]], 3
+// CHECK: [[TMP43:%.*]] = icmp eq i64 [[TMP42]], 0
+// CHECK: br i1 [[TMP43]], label [[OMP_TYPE_ALLOC9:%.*]], label [[OMP_TYPE_ALLOC_ELSE10:%.*]]
+// CHECK: omp.type.alloc9:
+// CHECK: [[TMP44:%.*]] = and i64 [[TMP41]], -4
+// CHECK: br label [[OMP_TYPE_END14:%.*]]
+// CHECK: omp.type.alloc.else10:
+// CHECK: [[TMP45:%.*]] = icmp eq i64 [[TMP42]], 1
+// CHECK: br i1 [[TMP45]], label [[OMP_TYPE_TO11:%.*]], label [[OMP_TYPE_TO_ELSE12:%.*]]
+// CHECK: omp.type.to11:
+// CHECK: [[TMP46:%.*]] = and i64 [[TMP41]], -3
+// CHECK: br label [[OMP_TYPE_END14]]
+// CHECK: omp.type.to.else12:
+// CHECK: [[TMP47:%.*]] = icmp eq i64 [[TMP42]], 2
+// CHECK: br i1 [[TMP47]], label [[OMP_TYPE_FROM13:%.*]], label [[OMP_TYPE_END14]]
+// CHECK: omp.type.from13:
+// CHECK: [[TMP48:%.*]] = and i64 [[TMP41]], -2
+// CHECK: br label [[OMP_TYPE_END14]]
+// CHECK: omp.type.end14:
+// CHECK: [[OMP_MAPTYPE15:%.*]] = phi i64 [ [[TMP44]], [[OMP_TYPE_ALLOC9]] ], [ [[TMP46]], [[OMP_TYPE_TO11]] ], [ [[TMP48]], [[OMP_TYPE_FROM13]] ], [ [[TMP41]], [[OMP_TYPE_TO_ELSE12]] ]
+// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP7]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[Z]], i64 4, i64 [[OMP_MAPTYPE15]], ptr null)
+// CHECK: [[TMP49:%.*]] = add nuw i64 281474976710675, [[TMP32]]
+// CHECK: [[TMP50:%.*]] = and i64 [[TMP12]], 3
+// CHECK: [[TMP51:%.*]] = icmp eq i64 [[TMP50]], 0
+// CHECK: br i1 [[TMP51]], label [[OMP_TYPE_ALLOC16:%.*]], label [[OMP_TYPE_ALLOC_ELSE17:%.*]]
+// CHECK: omp.type.alloc16:
+// CHECK: [[TMP52:%.*]] = and i64 [[TMP49]], -4
+// CHECK: br label [[OMP_TYPE_END21:%.*]]
+// CHECK: omp.type.alloc.else17:
+// CHECK: [[TMP53:%.*]] = icmp eq i64 [[TMP50]], 1
+// CHECK: br i1 [[TMP53]], label [[OMP_TYPE_TO18:%.*]], label [[OMP_TYPE_TO_ELSE19:%.*]]
+// CHECK: omp.type.to18:
+// CHECK: [[TMP54:%.*]] = and i64 [[TMP49]], -3
+// CHECK: br label [[OMP_TYPE_END21]]
+// CHECK: omp.type.to.else19:
+// CHECK: [[TMP55:%.*]] = icmp eq i64 [[TMP50]], 2
+// CHECK: br i1 [[TMP55]], label [[OMP_TYPE_FROM20:%.*]], label [[OMP_TYPE_END21]]
+// CHECK: omp.type.from20:
+// CHECK: [[TMP56:%.*]] = and i64 [[TMP49]], -2
+// CHECK: br label [[OMP_TYPE_END21]]
+// CHECK: omp.type.end21:
+// CHECK: [[OMP_MAPTYPE22:%.*]] = phi i64 [ [[TMP52]], [[OMP_TYPE_ALLOC16]] ], [ [[TMP54]], [[OMP_TYPE_TO18]] ], [ [[TMP56]], [[OMP_TYPE_FROM20]] ], [ [[TMP49]], [[OMP_TYPE_TO_ELSE19]] ]
+// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP7]], ptr [[S1P]], ptr [[X]], i64 4, i64 [[OMP_MAPTYPE22]], ptr null)
+// CHECK: [[TMP57:%.*]] = add nuw i64 281474976710675, [[TMP32]]
+// CHECK: [[TMP58:%.*]] = and i64 [[TMP12]], 3
+// CHECK: [[TMP59:%.*]] = icmp eq i64 [[TMP58]], 0
+// CHECK: br i1 [[TMP59]], label [[OMP_TYPE_ALLOC23:%.*]], label [[OMP_TYPE_ALLOC_ELSE24:%.*]]
+// CHECK: omp.type.alloc23:
+// CHECK: [[TMP60:%.*]] = and i64 [[TMP57]], -4
+// CHECK: br label [[OMP_TYPE_END28]]
+// CHECK: omp.type.alloc.else24:
+// CHECK: [[TMP61:%.*]] = icmp eq i64 [[TMP58]], 1
+// CHECK: br i1 [[TMP61]], label [[OMP_TYPE_TO25:%.*]], label [[OMP_TYPE_TO_ELSE26:%.*]]
+// CHECK: omp.type.to25:
+// CHECK: [[TMP62:%.*]] = and i64 [[TMP57]], -3
+// CHECK: br label [[OMP_TYPE_END28]]
+// CHECK: omp.type.to.else26:
+// CHECK: [[TMP63:%.*]] = icmp eq i64 [[TMP58]], 2
+// CHECK: br i1 [[TMP63]], label [[OMP_TYPE_FROM27:%.*]], label [[OMP_TYPE_END28]]
+// CHECK: omp.type.from27:
+// CHECK: [[TMP64:%.*]] = and i64 [[TMP57]], -2
+// CHECK: br label [[OMP_TYPE_END28]]
+// CHECK: omp.type.end28:
+// CHECK: [[OMP_MAPTYPE29:%.*]] = phi i64 [ [[TMP60]], [[OMP_TYPE_ALLOC23]] ], [ [[TMP62]], [[OMP_TYPE_TO25]] ], [ [[TMP64]], [[OMP_TYPE_FROM27]] ], [ [[TMP57]], [[OMP_TYPE_TO_ELSE26]] ]
+// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP7]], ptr [[S1P7]], ptr [[Y]], i64 4, i64 [[OMP_MAPTYPE29]], ptr null)
+// CHECK: [[OMP_ARRAYMAP_NEXT]] = getelementptr [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 1
+// CHECK: [[OMP_ARRAYMAP_ISDONE:%.*]] = icmp eq ptr [[OMP_ARRAYMAP_NEXT]], [[TMP11]]
+// CHECK: br i1 [[OMP_ARRAYMAP_ISDONE]], label [[OMP_ARRAYMAP_EXIT:%.*]], label [[OMP_ARRAYMAP_BODY]]
+// CHECK: omp.arraymap.exit:
+// CHECK: [[OMP_ARRAYINIT_ISARRAY30:%.*]] = icmp sgt i64 [[TMP10]], 1
+// CHECK: [[TMP65:%.*]] = and i64 [[TMP12]], 8
+// CHECK: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP65]], 0
+// CHECK: [[TMP66:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY30]], [[DOTOMP_ARRAY__DEL__DELETE]]
+// CHECK: br i1 [[TMP66]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]]
+// CHECK: .omp.array..del:
+// CHECK: [[TMP67:%.*]] = mul nuw i64 [[TMP10]], 16
+// CHECK: [[TMP68:%.*]] = and i64 [[TMP12]], -4
+// CHECK: [[TMP69:%.*]] = or i64 [[TMP68]], 512
+// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP7]], ptr [[TMP8]], ptr [[TMP9]], i64 [[TMP67]], i64 [[TMP69]], ptr [[TMP13]])
+// CHECK: br label [[OMP_DONE]]
+// CHECK: omp.done:
+// CHECK: ret void
+//
>From d1ea96f088880a9dd22670b6c09752fb0910b136 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Fri, 24 Apr 2026 03:35:25 -0700
Subject: [PATCH 08/13] Update more clang/offload tests.
---
clang/test/OpenMP/declare_mapper_codegen.cpp | 142 ++++----
..._of_structs_with_nested_mapper_codegen.cpp | 164 +++++-----
..._of_structs_with_nested_mapper_codegen.cpp | 164 +++++-----
...t_map_nested_ptr_member_mapper_codegen.cpp | 304 +++++++++---------
.../mapping/declare_mapper_nested_mappers.cpp | 7 +-
.../test/mapping/declare_mapper_target.cpp | 2 +-
.../mapping/declare_mapper_target_data.cpp | 2 +-
.../declare_mapper_target_data_enter_exit.cpp | 2 +-
.../mapping/declare_mapper_target_update.cpp | 2 +-
9 files changed, 426 insertions(+), 363 deletions(-)
diff --git a/clang/test/OpenMP/declare_mapper_codegen.cpp b/clang/test/OpenMP/declare_mapper_codegen.cpp
index ff891396fe72e..a67cb9dc246ab 100644
--- a/clang/test/OpenMP/declare_mapper_codegen.cpp
+++ b/clang/test/OpenMP/declare_mapper_codegen.cpp
@@ -85,6 +85,11 @@ class C {
};
#pragma omp declare mapper(id: C s) map(s.a, s.b[0:2])
+//
+// Per-element entries (N = __tgt_mapper_num_components()):
+// &s, &s.a, sizeof(int), MEMBER_OF(N) | TO | FROM | modifiers
+// &s.b[0], &s.b[0], sizeof(double)*2, TO | FROM | modifiers
+// &s.b, &s.b[0], sizeof(double*), ATTACH
// CK0: define {{.*}}void [[MPRFUNC:@[.]omp_mapper[.].*C[.]id]](ptr noundef [[HANDLE:%.+]], ptr noundef [[BPTR:%.+]], ptr noundef [[BEGIN:%.+]], i64 noundef [[BYTESIZE:%.+]], i64 noundef [[TYPE:%.+]], ptr{{.*}})
// CK0-64-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], 16
@@ -114,20 +119,14 @@ class C {
// CK0: [[PTR:%.+]] = phi ptr [ [[BEGIN]], %{{.+}} ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
// CK0-DAG: [[ABEGIN:%.+]] = getelementptr inbounds nuw %class.C, ptr [[PTR]], i32 0, i32 0
// CK0-DAG: [[BBEGIN:%.+]] = getelementptr inbounds nuw %class.C, ptr [[PTR]], i32 0, i32 1
+// CK0-DAG: [[BARRBASE:%.+]] = load ptr, ptr [[BBEGIN]]
// CK0-DAG: [[BBEGIN2:%.+]] = getelementptr inbounds nuw %class.C, ptr [[PTR]], i32 0, i32 1
// CK0-DAG: [[BARRBEGIN:%.+]] = load ptr, ptr [[BBEGIN2]]
// CK0-DAG: [[BARRBEGINGEP:%.+]] = getelementptr inbounds nuw double, ptr [[BARRBEGIN]], i[[sz:64|32]] 0
-// CK0-DAG: [[BEND:%.+]] = getelementptr ptr, ptr [[BBEGIN]], i32 1
-// CK0-64-DAG: [[ABEGINI:%.+]] = ptrtoaddr ptr [[ABEGIN]] to i64
-// CK0-64-DAG: [[BENDI:%.+]] = ptrtoaddr ptr [[BEND]] to i64
-// CK0-64-DAG: [[CUSIZE:%.+]] = sub i64 [[BENDI]], [[ABEGINI]]
-// CK0-32-DAG: [[ABEGINI:%.+]] = ptrtoaddr ptr [[ABEGIN]] to i32
-// CK0-32-DAG: [[BENDI:%.+]] = ptrtoaddr ptr [[BEND]] to i32
-// CK0-32-DAG: [[CSIZE:%.+]] = sub i32 [[BENDI]], [[ABEGINI]]
-// CK0-32-DAG: [[CUSIZE:%.+]] = zext i32 [[CSIZE]] to i64
// CK0-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(ptr [[HANDLE]])
// CK0-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 48
-// CK0-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 0, [[SHIPRESIZE]]
+// &s, &s.a, sizeof(int), MEMBER_OF(N) | TO | FROM | modifiers
+// CK0-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 3, [[SHIPRESIZE]]
// CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
// CK0-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
// CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
@@ -148,53 +147,50 @@ class C {
// CK0-DAG: br label %[[TYEND]]
// CK0-DAG: [[TYEND]]
// CK0-DAG: [[PHITYPE0:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
-// CK0: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 [[CUSIZE]], i64 [[PHITYPE0]], {{.*}})
-// 281474976710659 == 0x1,000,000,003
-// CK0-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710659, [[SHIPRESIZE]]
+// CK0-DAG: [[MODMASK0:%.+]] = and i64 [[TYPE]], 5132
+// CK0-DAG: [[PHITYPE0_MOD:%.+]] = or i64 [[PHITYPE0]], [[MODMASK0]]
+// CK0: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 4, i64 [[PHITYPE0_MOD]], {{.*}})
+// &s.b[0], &s.b[0], sizeof(double)*2, TO | FROM | modifiers
// CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
// CK0-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
// CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
// CK0-DAG: [[ALLOC]]
-// CK0-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
// CK0-DAG: br label %[[TYEND:[^,]+]]
// CK0-DAG: [[ALLOCELSE]]
// CK0-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
// CK0-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
// CK0-DAG: [[TO]]
-// CK0-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
// CK0-DAG: br label %[[TYEND]]
// CK0-DAG: [[TOELSE]]
// CK0-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
// CK0-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
// CK0-DAG: [[FROM]]
-// CK0-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
// CK0-DAG: br label %[[TYEND]]
// CK0-DAG: [[TYEND]]
-// CK0-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
-// CK0: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 4, i64 [[TYPE1]], {{.*}})
-// 281474976710675 == 0x1,000,000,013
-// CK0-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710675, [[SHIPRESIZE]]
+// CK0-DAG: [[TYPE1:%.+]] = phi i64 [ 0, %[[ALLOC]] ], [ 1, %[[TO]] ], [ 2, %[[FROM]] ], [ 3, %[[TOELSE]] ]
+// CK0-DAG: [[MODMASK1:%.+]] = and i64 [[TYPE]], 5132
+// CK0-DAG: [[TYPE1_MOD:%.+]] = or i64 [[TYPE1]], [[MODMASK1]]
+// CK0: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BARRBASE]], ptr [[BARRBEGINGEP]], i64 16, i64 [[TYPE1_MOD]], {{.*}})
+// &s.b, &s.b[0], sizeof(double*), ATTACH
// CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
// CK0-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
// CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
// CK0-DAG: [[ALLOC]]
-// CK0-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
// CK0-DAG: br label %[[TYEND:[^,]+]]
// CK0-DAG: [[ALLOCELSE]]
// CK0-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
// CK0-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
// CK0-DAG: [[TO]]
-// CK0-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
// CK0-DAG: br label %[[TYEND]]
// CK0-DAG: [[TOELSE]]
// CK0-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
// CK0-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
// CK0-DAG: [[FROM]]
-// CK0-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
// CK0-DAG: br label %[[TYEND]]
// CK0-DAG: [[TYEND]]
-// CK0-DAG: [[TYPE2:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
-// CK0: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BBEGIN]], ptr [[BARRBEGINGEP]], i64 16, i64 [[TYPE2]], {{.*}})
+// CK0-DAG: [[TYPE2:%.+]] = phi i64 [ 16384, %[[ALLOC]] ], [ 16384, %[[TO]] ], [ 16384, %[[FROM]] ], [ 16384, %[[TOELSE]] ]
+// CK0-64: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BBEGIN]], ptr [[BARRBEGINGEP]], i64 8, i64 [[TYPE2]], {{.*}})
+// CK0-32: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BBEGIN]], ptr [[BARRBEGINGEP]], i64 4, i64 [[TYPE2]], {{.*}})
// CK0: [[PTRNEXT]] = getelementptr %class.C, ptr [[PTR]], i32 1
// CK0: [[ISDONE:%.+]] = icmp eq ptr [[PTRNEXT]], [[PTREND]]
// CK0: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
@@ -586,6 +582,9 @@ class C {
};
#pragma omp declare mapper(id: C<int> s) map(s.a)
+//
+// Per-element entries (N = __tgt_mapper_num_components()):
+// &s, &s.a, sizeof(int), MEMBER_OF(N) | TO | FROM | modifiers
// CK1: define {{.*}}void @.omp_mapper.{{.*}}C{{.*}}.id{{.*}}(ptr noundef [[HANDLE:%.+]], ptr noundef [[BPTR:%.+]], ptr noundef [[BEGIN:%.+]], i64 noundef [[BYTESIZE:%.+]], i64 noundef [[TYPE:%.+]], ptr{{.*}})
// CK1-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], 4
@@ -636,7 +635,9 @@ class C {
// CK1-DAG: br label %[[TYEND]]
// CK1-DAG: [[TYEND]]
// CK1-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
-// CK1: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 4, i64 [[TYPE1]], {{.*}})
+// CK1-DAG: [[MODMASK:%.+]] = and i64 [[TYPE]], 5132
+// CK1-DAG: [[TYPE1_MOD:%.+]] = or i64 [[TYPE1]], [[MODMASK]]
+// CK1: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 4, i64 [[TYPE1_MOD]], {{.*}})
// CK1: [[PTRNEXT]] = getelementptr %class.C, ptr [[PTR]], i32 1
// CK1: [[ISDONE:%.+]] = icmp eq ptr [[PTRNEXT]], [[PTREND]]
// CK1: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
@@ -691,6 +692,9 @@ class C {
#pragma omp declare mapper(B s) map(s.a)
#pragma omp declare mapper(id: C s) map(s.b)
+//
+// Per-element entries emitted (N = __tgt_mapper_num_components()):
+// &s, &s.b, sizeof(B), MEMBER_OF(N) | TO | FROM | modifiers (dispatches to B mapper)
// CK2: define {{.*}}void [[BMPRFUNC:@[.]omp_mapper[.].*B[.]default]](ptr{{.*}}, ptr{{.*}}, ptr{{.*}}, i64{{.*}}, i64{{.*}}, ptr{{.*}})
@@ -743,7 +747,9 @@ class C {
// CK2-DAG: br label %[[TYEND]]
// CK2-DAG: [[TYEND]]
// CK2-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
-// CK2: call void [[BMPRFUNC]](ptr [[HANDLE]], ptr [[PTR]], ptr [[BBEGIN]], i64 8, i64 [[TYPE1]], {{.*}})
+// CK2-DAG: [[MODMASK:%.+]] = and i64 [[TYPE]], 5132
+// CK2-DAG: [[TYPE1_MOD:%.+]] = or i64 [[TYPE1]], [[MODMASK]]
+// CK2: call void [[BMPRFUNC]](ptr [[HANDLE]], ptr [[PTR]], ptr [[BBEGIN]], i64 8, i64 [[TYPE1_MOD]], {{.*}})
// CK2: [[PTRNEXT]] = getelementptr %class.C, ptr [[PTR]], i32 1
// CK2: [[ISDONE:%.+]] = icmp eq ptr [[PTRNEXT]], [[PTREND]]
// CK2: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
@@ -884,6 +890,11 @@ class C {
};
#pragma omp declare mapper(id: C s) map(s.a, s.b[0:2])
+//
+// Per-element entries (N = __tgt_mapper_num_components()):
+// &s, &s.a, sizeof(int), MEMBER_OF(N) | TO | FROM | modifiers
+// &s.b[0], &s.b[0], sizeof(double)*2, TO | FROM | modifiers
+// &s.b, &s.b[0], sizeof(double*), ATTACH
// CK4: define {{.*}}void [[MPRFUNC:@[.]omp_mapper[.].*C[.]id]](ptr noundef [[HANDLE:%.+]], ptr noundef [[BPTR:%.+]], ptr noundef [[BEGIN:%.+]], i64 noundef [[BYTESIZE:%.+]], i64 noundef [[TYPE:%.+]], ptr{{.*}})
// CK4-64-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], 16
@@ -914,20 +925,14 @@ class C {
// CK4: [[PTR:%.+]] = phi ptr [ [[BEGIN]], %{{.+}} ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
// CK4-DAG: [[ABEGIN:%.+]] = getelementptr inbounds nuw %class.C, ptr [[PTR]], i32 0, i32 0
// CK4-DAG: [[BBEGIN:%.+]] = getelementptr inbounds nuw %class.C, ptr [[PTR]], i32 0, i32 1
+// CK4-DAG: [[BARRBASE:%.+]] = load ptr, ptr [[BBEGIN]]
// CK4-DAG: [[BBEGIN2:%.+]] = getelementptr inbounds nuw %class.C, ptr [[PTR]], i32 0, i32 1
// CK4-DAG: [[BARRBEGIN:%.+]] = load ptr, ptr [[BBEGIN2]]
// CK4-DAG: [[BARRBEGINGEP:%.+]] = getelementptr inbounds nuw double, ptr [[BARRBEGIN]], i[[sz:64|32]] 0
-// CK4-DAG: [[BEND:%.+]] = getelementptr ptr, ptr [[BBEGIN]], i32 1
-// CK4-64-DAG: [[ABEGINI:%.+]] = ptrtoaddr ptr [[ABEGIN]] to i64
-// CK4-64-DAG: [[BENDI:%.+]] = ptrtoaddr ptr [[BEND]] to i64
-// CK4-64-DAG: [[CUSIZE:%.+]] = sub i64 [[BENDI]], [[ABEGINI]]
-// CK4-32-DAG: [[ABEGINI:%.+]] = ptrtoaddr ptr [[ABEGIN]] to i32
-// CK4-32-DAG: [[BENDI:%.+]] = ptrtoaddr ptr [[BEND]] to i32
-// CK4-32-DAG: [[CSIZE:%.+]] = sub i32 [[BENDI]], [[ABEGINI]]
-// CK4-32-DAG: [[CUSIZE:%.+]] = zext i32 [[CSIZE]] to i64
// CK4-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(ptr [[HANDLE]])
// CK4-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 48
-// CK4-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 0, [[SHIPRESIZE]]
+// &s, &s.a, sizeof(int), MEMBER_OF(N) | TO | FROM | modifiers
+// CK4-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 3, [[SHIPRESIZE]]
// CK4-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
// CK4-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
// CK4-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
@@ -948,53 +953,50 @@ class C {
// CK4-DAG: br label %[[TYEND]]
// CK4-DAG: [[TYEND]]
// CK4-DAG: [[PHITYPE0:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
-// CK4: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 [[CUSIZE]], i64 [[PHITYPE0]], {{.*}})
-// 281474976710659 == 0x1,000,000,003
-// CK4-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710659, [[SHIPRESIZE]]
+// CK4-DAG: [[MODMASK0:%.+]] = and i64 [[TYPE]], 5132
+// CK4-DAG: [[PHITYPE0_MOD:%.+]] = or i64 [[PHITYPE0]], [[MODMASK0]]
+// CK4: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 4, i64 [[PHITYPE0_MOD]], {{.*}})
+// &s.b[0], &s.b[0], sizeof(double)*2, TO | FROM | modifiers
// CK4-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
// CK4-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
// CK4-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
// CK4-DAG: [[ALLOC]]
-// CK4-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
// CK4-DAG: br label %[[TYEND:[^,]+]]
// CK4-DAG: [[ALLOCELSE]]
// CK4-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
// CK4-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
// CK4-DAG: [[TO]]
-// CK4-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
// CK4-DAG: br label %[[TYEND]]
// CK4-DAG: [[TOELSE]]
// CK4-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
// CK4-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
// CK4-DAG: [[FROM]]
-// CK4-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
// CK4-DAG: br label %[[TYEND]]
// CK4-DAG: [[TYEND]]
-// CK4-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
-// CK4: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 4, i64 [[TYPE1]], {{.*}})
-// 281474976710675 == 0x1,000,000,013
-// CK4-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710675, [[SHIPRESIZE]]
+// CK4-DAG: [[TYPE1:%.+]] = phi i64 [ 0, %[[ALLOC]] ], [ 1, %[[TO]] ], [ 2, %[[FROM]] ], [ 3, %[[TOELSE]] ]
+// CK4-DAG: [[MODMASK1:%.+]] = and i64 [[TYPE]], 5132
+// CK4-DAG: [[TYPE1_MOD:%.+]] = or i64 [[TYPE1]], [[MODMASK1]]
+// CK4: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BARRBASE]], ptr [[BARRBEGINGEP]], i64 16, i64 [[TYPE1_MOD]], {{.*}})
+// &s.b, &s.b[0], sizeof(double*), ATTACH
// CK4-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
// CK4-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
// CK4-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
// CK4-DAG: [[ALLOC]]
-// CK4-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
// CK4-DAG: br label %[[TYEND:[^,]+]]
// CK4-DAG: [[ALLOCELSE]]
// CK4-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
// CK4-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
// CK4-DAG: [[TO]]
-// CK4-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
// CK4-DAG: br label %[[TYEND]]
// CK4-DAG: [[TOELSE]]
// CK4-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
// CK4-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
// CK4-DAG: [[FROM]]
-// CK4-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
// CK4-DAG: br label %[[TYEND]]
// CK4-DAG: [[TYEND]]
-// CK4-DAG: [[TYPE2:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
-// CK4: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BBEGIN]], ptr [[BARRBEGINGEP]], i64 16, i64 [[TYPE2]], {{.*}})
+// CK4-DAG: [[TYPE2:%.+]] = phi i64 [ 16384, %[[ALLOC]] ], [ 16384, %[[TO]] ], [ 16384, %[[FROM]] ], [ 16384, %[[TOELSE]] ]
+// CK4-64: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BBEGIN]], ptr [[BARRBEGINGEP]], i64 8, i64 [[TYPE2]], {{.*}})
+// CK4-32: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BBEGIN]], ptr [[BARRBEGINGEP]], i64 4, i64 [[TYPE2]], {{.*}})
// CK4: [[PTRNEXT]] = getelementptr %class.C, ptr [[PTR]], i32 1
// CK4: [[ISDONE:%.+]] = icmp eq ptr [[PTRNEXT]], [[PTREND]]
// CK4: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
@@ -1067,9 +1069,14 @@ typedef struct myvec {
} myvec_t;
#pragma omp declare mapper(id: myvec_t v) map(iterator(it=0:v.a), tofrom: v.b[it])
+//
+// Per-element entries emitted for struct v (N = __tgt_mapper_num_components()):
+// &v.b[it], &v.b[it], sizeof(double), TO | FROM | modifiers
+// &v.b, &v.b[it], sizeof(double*), ATTACH
+
// CK5: @[[ITER:[a-zA-Z0-9_]+]] = global i32 0, align 4
-void foo(){
+void foo(){
myvec_t s;
#pragma omp target map(mapper(id), to:s)
{
@@ -1101,32 +1108,51 @@ void foo(){
// CK5: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]]
// CK5: [[LBODY]]
// CK5: [[PTR:%.+]] = phi ptr [ [[BEGIN]], %{{.+}} ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
-// CK5-DAG: [[ABEGIN:%.+]] = getelementptr inbounds nuw %struct.myvec, ptr [[PTR]], i32 0, i32 1
+// CK5-DAG: [[BBEGIN:%.+]] = getelementptr inbounds nuw %struct.myvec, ptr [[PTR]], i32 0, i32 1
+// CK5-DAG: [[BBASE:%.+]] = load ptr, ptr [[BBEGIN]], align {{.*}}
// CK5-DAG: load i32, ptr @[[ITER]], align 4
// CK5-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(ptr [[HANDLE]])
// CK5-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 48
-// CK5-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 0, [[SHIPRESIZE]]
+// &v.b[it], &v.b[it], sizeof(double), TO | FROM | modifiers
+// CK5-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
+// CK5-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
+// CK5-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
+// CK5-DAG: [[ALLOC]]
+// CK5-DAG: br label %[[TYEND:[^,]+]]
+// CK5-DAG: [[ALLOCELSE]]
+// CK5-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
+// CK5-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
+// CK5-DAG: [[TO]]
+// CK5-DAG: br label %[[TYEND]]
+// CK5-DAG: [[TOELSE]]
+// CK5-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
+// CK5-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
+// CK5-DAG: [[FROM]]
+// CK5-DAG: br label %[[TYEND]]
+// CK5-DAG: [[TYEND]]
+// CK5-DAG: [[TYPE1:%.+]] = phi i64 [ 0, %[[ALLOC]] ], [ 1, %[[TO]] ], [ 2, %[[FROM]] ], [ 3, %[[TOELSE]] ]
+// CK5-DAG: [[MODMASK:%.+]] = and i64 [[TYPE]], 5132
+// CK5-DAG: [[TYPE1_MOD:%.+]] = or i64 [[TYPE1]], [[MODMASK]]
+// CK5: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BBASE]], ptr {{.*}}, i64 {{.*}}, i64 [[TYPE1_MOD]], {{.*}})
+// &v.b, &v.b[it], sizeof(double*), ATTACH
// CK5-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
// CK5-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
// CK5-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
// CK5-DAG: [[ALLOC]]
-// CK5-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
// CK5-DAG: br label %[[TYEND:[^,]+]]
// CK5-DAG: [[ALLOCELSE]]
// CK5-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
// CK5-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
// CK5-DAG: [[TO]]
-// CK5-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
// CK5-DAG: br label %[[TYEND]]
// CK5-DAG: [[TOELSE]]
// CK5-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
// CK5-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
// CK5-DAG: [[FROM]]
-// CK5-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
// CK5-DAG: br label %[[TYEND]]
// CK5-DAG: [[TYEND]]
-// CK5-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
-// CK5: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 {{.*}}, i64 [[TYPE1]], {{.*}})
+// CK5-DAG: [[TYPE2:%.+]] = phi i64 [ 16384, %[[ALLOC]] ], [ 16384, %[[TO]] ], [ 16384, %[[FROM]] ], [ 16384, %[[TOELSE]] ]
+// CK5: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BBEGIN]], ptr {{.*}}, i64 {{.*}}, i64 [[TYPE2]], {{.*}})
// CK5: [[PTRNEXT]] = getelementptr %struct.myvec, ptr [[PTR]], i32 1
// CK5: [[ISDONE:%.+]] = icmp eq ptr [[PTRNEXT]], [[PTREND]]
// CK5: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
diff --git a/clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_codegen.cpp b/clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_codegen.cpp
index 1b9fbb9868916..5703205167c8f 100644
--- a/clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_codegen.cpp
+++ b/clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_codegen.cpp
@@ -139,7 +139,7 @@ void foo() {
// CHECK-NEXT: [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP2]], [[TMP7]]
// CHECK-NEXT: br i1 [[OMP_ARRAYMAP_ISEMPTY]], label [[OMP_DONE:%.*]], label [[OMP_ARRAYMAP_BODY:%.*]]
// CHECK: omp.arraymap.body:
-// CHECK-NEXT: [[OMP_ARRAYMAP_PTRCURRENT:%.*]] = phi ptr [ [[TMP2]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END20:%.*]] ]
+// CHECK-NEXT: [[OMP_ARRAYMAP_PTRCURRENT:%.*]] = phi ptr [ [[TMP2]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END22:%.*]] ]
// CHECK-NEXT: [[E:%.*]] = getelementptr inbounds nuw [[STRUCT_D]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
// CHECK-NEXT: [[F:%.*]] = getelementptr inbounds nuw [[STRUCT_D]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 1
// CHECK-NEXT: [[H:%.*]] = getelementptr inbounds nuw [[STRUCT_D]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 2
@@ -170,87 +170,95 @@ void foo() {
// CHECK-NEXT: br label [[OMP_TYPE_END]]
// CHECK: omp.type.end:
// CHECK-NEXT: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP24]], [[OMP_TYPE_ALLOC]] ], [ [[TMP26]], [[OMP_TYPE_TO]] ], [ [[TMP28]], [[OMP_TYPE_FROM]] ], [ [[TMP21]], [[OMP_TYPE_TO_ELSE]] ]
-// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 [[TMP18]], i64 [[OMP_MAPTYPE]], ptr null)
-// CHECK-NEXT: [[TMP29:%.*]] = add nuw i64 281474976711171, [[TMP20]]
-// CHECK-NEXT: [[TMP30:%.*]] = and i64 [[TMP4]], 3
-// CHECK-NEXT: [[TMP31:%.*]] = icmp eq i64 [[TMP30]], 0
-// CHECK-NEXT: br i1 [[TMP31]], label [[OMP_TYPE_ALLOC1:%.*]], label [[OMP_TYPE_ALLOC_ELSE2:%.*]]
+// CHECK-NEXT: [[TMP29:%.*]] = and i64 [[TMP4]], 5132
+// CHECK-NEXT: [[OMP_MAPTYPE_WITH_MODIFIERS:%.*]] = or i64 [[OMP_MAPTYPE]], [[TMP29]]
+// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 [[TMP18]], i64 [[OMP_MAPTYPE_WITH_MODIFIERS]], ptr null)
+// CHECK-NEXT: [[TMP30:%.*]] = add nuw i64 281474976711171, [[TMP20]]
+// CHECK-NEXT: [[TMP31:%.*]] = and i64 [[TMP4]], 3
+// CHECK-NEXT: [[TMP32:%.*]] = icmp eq i64 [[TMP31]], 0
+// CHECK-NEXT: br i1 [[TMP32]], label [[OMP_TYPE_ALLOC1:%.*]], label [[OMP_TYPE_ALLOC_ELSE2:%.*]]
// CHECK: omp.type.alloc1:
-// CHECK-NEXT: [[TMP32:%.*]] = and i64 [[TMP29]], -4
+// CHECK-NEXT: [[TMP33:%.*]] = and i64 [[TMP30]], -4
// CHECK-NEXT: br label [[OMP_TYPE_END6:%.*]]
// CHECK: omp.type.alloc.else2:
-// CHECK-NEXT: [[TMP33:%.*]] = icmp eq i64 [[TMP30]], 1
-// CHECK-NEXT: br i1 [[TMP33]], label [[OMP_TYPE_TO3:%.*]], label [[OMP_TYPE_TO_ELSE4:%.*]]
+// CHECK-NEXT: [[TMP34:%.*]] = icmp eq i64 [[TMP31]], 1
+// CHECK-NEXT: br i1 [[TMP34]], label [[OMP_TYPE_TO3:%.*]], label [[OMP_TYPE_TO_ELSE4:%.*]]
// CHECK: omp.type.to3:
-// CHECK-NEXT: [[TMP34:%.*]] = and i64 [[TMP29]], -3
+// CHECK-NEXT: [[TMP35:%.*]] = and i64 [[TMP30]], -3
// CHECK-NEXT: br label [[OMP_TYPE_END6]]
// CHECK: omp.type.to.else4:
-// CHECK-NEXT: [[TMP35:%.*]] = icmp eq i64 [[TMP30]], 2
-// CHECK-NEXT: br i1 [[TMP35]], label [[OMP_TYPE_FROM5:%.*]], label [[OMP_TYPE_END6]]
+// CHECK-NEXT: [[TMP36:%.*]] = icmp eq i64 [[TMP31]], 2
+// CHECK-NEXT: br i1 [[TMP36]], label [[OMP_TYPE_FROM5:%.*]], label [[OMP_TYPE_END6]]
// CHECK: omp.type.from5:
-// CHECK-NEXT: [[TMP36:%.*]] = and i64 [[TMP29]], -2
+// CHECK-NEXT: [[TMP37:%.*]] = and i64 [[TMP30]], -2
// CHECK-NEXT: br label [[OMP_TYPE_END6]]
// CHECK: omp.type.end6:
-// CHECK-NEXT: [[OMP_MAPTYPE7:%.*]] = phi i64 [ [[TMP32]], [[OMP_TYPE_ALLOC1]] ], [ [[TMP34]], [[OMP_TYPE_TO3]] ], [ [[TMP36]], [[OMP_TYPE_FROM5]] ], [ [[TMP29]], [[OMP_TYPE_TO_ELSE4]] ]
-// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 4, i64 [[OMP_MAPTYPE7]], ptr null)
-// CHECK-NEXT: [[TMP37:%.*]] = add nuw i64 281474976711171, [[TMP20]]
-// CHECK-NEXT: [[TMP38:%.*]] = and i64 [[TMP4]], 3
-// CHECK-NEXT: [[TMP39:%.*]] = icmp eq i64 [[TMP38]], 0
-// CHECK-NEXT: br i1 [[TMP39]], label [[OMP_TYPE_ALLOC8:%.*]], label [[OMP_TYPE_ALLOC_ELSE9:%.*]]
-// CHECK: omp.type.alloc8:
-// CHECK-NEXT: [[TMP40:%.*]] = and i64 [[TMP37]], -4
-// CHECK-NEXT: br label [[OMP_TYPE_END13:%.*]]
-// CHECK: omp.type.alloc.else9:
-// CHECK-NEXT: [[TMP41:%.*]] = icmp eq i64 [[TMP38]], 1
-// CHECK-NEXT: br i1 [[TMP41]], label [[OMP_TYPE_TO10:%.*]], label [[OMP_TYPE_TO_ELSE11:%.*]]
-// CHECK: omp.type.to10:
-// CHECK-NEXT: [[TMP42:%.*]] = and i64 [[TMP37]], -3
-// CHECK-NEXT: br label [[OMP_TYPE_END13]]
-// CHECK: omp.type.to.else11:
-// CHECK-NEXT: [[TMP43:%.*]] = icmp eq i64 [[TMP38]], 2
-// CHECK-NEXT: br i1 [[TMP43]], label [[OMP_TYPE_FROM12:%.*]], label [[OMP_TYPE_END13]]
-// CHECK: omp.type.from12:
-// CHECK-NEXT: [[TMP44:%.*]] = and i64 [[TMP37]], -2
-// CHECK-NEXT: br label [[OMP_TYPE_END13]]
-// CHECK: omp.type.end13:
-// CHECK-NEXT: [[OMP_MAPTYPE14:%.*]] = phi i64 [ [[TMP40]], [[OMP_TYPE_ALLOC8]] ], [ [[TMP42]], [[OMP_TYPE_TO10]] ], [ [[TMP44]], [[OMP_TYPE_FROM12]] ], [ [[TMP37]], [[OMP_TYPE_TO_ELSE11]] ]
-// CHECK-NEXT: call void @.omp_mapper._ZTS1C.default(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[F]], i64 4, i64 [[OMP_MAPTYPE14]], ptr null) #[[ATTR3]]
-// CHECK-NEXT: [[TMP45:%.*]] = add nuw i64 281474976711171, [[TMP20]]
-// CHECK-NEXT: [[TMP46:%.*]] = and i64 [[TMP4]], 3
-// CHECK-NEXT: [[TMP47:%.*]] = icmp eq i64 [[TMP46]], 0
-// CHECK-NEXT: br i1 [[TMP47]], label [[OMP_TYPE_ALLOC15:%.*]], label [[OMP_TYPE_ALLOC_ELSE16:%.*]]
-// CHECK: omp.type.alloc15:
-// CHECK-NEXT: [[TMP48:%.*]] = and i64 [[TMP45]], -4
-// CHECK-NEXT: br label [[OMP_TYPE_END20]]
-// CHECK: omp.type.alloc.else16:
-// CHECK-NEXT: [[TMP49:%.*]] = icmp eq i64 [[TMP46]], 1
-// CHECK-NEXT: br i1 [[TMP49]], label [[OMP_TYPE_TO17:%.*]], label [[OMP_TYPE_TO_ELSE18:%.*]]
-// CHECK: omp.type.to17:
-// CHECK-NEXT: [[TMP50:%.*]] = and i64 [[TMP45]], -3
-// CHECK-NEXT: br label [[OMP_TYPE_END20]]
-// CHECK: omp.type.to.else18:
-// CHECK-NEXT: [[TMP51:%.*]] = icmp eq i64 [[TMP46]], 2
-// CHECK-NEXT: br i1 [[TMP51]], label [[OMP_TYPE_FROM19:%.*]], label [[OMP_TYPE_END20]]
-// CHECK: omp.type.from19:
-// CHECK-NEXT: [[TMP52:%.*]] = and i64 [[TMP45]], -2
-// CHECK-NEXT: br label [[OMP_TYPE_END20]]
-// CHECK: omp.type.end20:
-// CHECK-NEXT: [[OMP_MAPTYPE21:%.*]] = phi i64 [ [[TMP48]], [[OMP_TYPE_ALLOC15]] ], [ [[TMP50]], [[OMP_TYPE_TO17]] ], [ [[TMP52]], [[OMP_TYPE_FROM19]] ], [ [[TMP45]], [[OMP_TYPE_TO_ELSE18]] ]
-// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[H]], i64 4, i64 [[OMP_MAPTYPE21]], ptr null)
+// CHECK-NEXT: [[OMP_MAPTYPE7:%.*]] = phi i64 [ [[TMP33]], [[OMP_TYPE_ALLOC1]] ], [ [[TMP35]], [[OMP_TYPE_TO3]] ], [ [[TMP37]], [[OMP_TYPE_FROM5]] ], [ [[TMP30]], [[OMP_TYPE_TO_ELSE4]] ]
+// CHECK-NEXT: [[TMP38:%.*]] = and i64 [[TMP4]], 5132
+// CHECK-NEXT: [[OMP_MAPTYPE_WITH_MODIFIERS8:%.*]] = or i64 [[OMP_MAPTYPE7]], [[TMP38]]
+// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS8]], ptr null)
+// CHECK-NEXT: [[TMP39:%.*]] = add nuw i64 281474976711171, [[TMP20]]
+// CHECK-NEXT: [[TMP40:%.*]] = and i64 [[TMP4]], 3
+// CHECK-NEXT: [[TMP41:%.*]] = icmp eq i64 [[TMP40]], 0
+// CHECK-NEXT: br i1 [[TMP41]], label [[OMP_TYPE_ALLOC9:%.*]], label [[OMP_TYPE_ALLOC_ELSE10:%.*]]
+// CHECK: omp.type.alloc9:
+// CHECK-NEXT: [[TMP42:%.*]] = and i64 [[TMP39]], -4
+// CHECK-NEXT: br label [[OMP_TYPE_END14:%.*]]
+// CHECK: omp.type.alloc.else10:
+// CHECK-NEXT: [[TMP43:%.*]] = icmp eq i64 [[TMP40]], 1
+// CHECK-NEXT: br i1 [[TMP43]], label [[OMP_TYPE_TO11:%.*]], label [[OMP_TYPE_TO_ELSE12:%.*]]
+// CHECK: omp.type.to11:
+// CHECK-NEXT: [[TMP44:%.*]] = and i64 [[TMP39]], -3
+// CHECK-NEXT: br label [[OMP_TYPE_END14]]
+// CHECK: omp.type.to.else12:
+// CHECK-NEXT: [[TMP45:%.*]] = icmp eq i64 [[TMP40]], 2
+// CHECK-NEXT: br i1 [[TMP45]], label [[OMP_TYPE_FROM13:%.*]], label [[OMP_TYPE_END14]]
+// CHECK: omp.type.from13:
+// CHECK-NEXT: [[TMP46:%.*]] = and i64 [[TMP39]], -2
+// CHECK-NEXT: br label [[OMP_TYPE_END14]]
+// CHECK: omp.type.end14:
+// CHECK-NEXT: [[OMP_MAPTYPE15:%.*]] = phi i64 [ [[TMP42]], [[OMP_TYPE_ALLOC9]] ], [ [[TMP44]], [[OMP_TYPE_TO11]] ], [ [[TMP46]], [[OMP_TYPE_FROM13]] ], [ [[TMP39]], [[OMP_TYPE_TO_ELSE12]] ]
+// CHECK-NEXT: [[TMP47:%.*]] = and i64 [[TMP4]], 5132
+// CHECK-NEXT: [[OMP_MAPTYPE_WITH_MODIFIERS16:%.*]] = or i64 [[OMP_MAPTYPE15]], [[TMP47]]
+// CHECK-NEXT: call void @.omp_mapper._ZTS1C.default(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[F]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS16]], ptr null) #[[ATTR3]]
+// CHECK-NEXT: [[TMP48:%.*]] = add nuw i64 281474976711171, [[TMP20]]
+// CHECK-NEXT: [[TMP49:%.*]] = and i64 [[TMP4]], 3
+// CHECK-NEXT: [[TMP50:%.*]] = icmp eq i64 [[TMP49]], 0
+// CHECK-NEXT: br i1 [[TMP50]], label [[OMP_TYPE_ALLOC17:%.*]], label [[OMP_TYPE_ALLOC_ELSE18:%.*]]
+// CHECK: omp.type.alloc17:
+// CHECK-NEXT: [[TMP51:%.*]] = and i64 [[TMP48]], -4
+// CHECK-NEXT: br label [[OMP_TYPE_END22]]
+// CHECK: omp.type.alloc.else18:
+// CHECK-NEXT: [[TMP52:%.*]] = icmp eq i64 [[TMP49]], 1
+// CHECK-NEXT: br i1 [[TMP52]], label [[OMP_TYPE_TO19:%.*]], label [[OMP_TYPE_TO_ELSE20:%.*]]
+// CHECK: omp.type.to19:
+// CHECK-NEXT: [[TMP53:%.*]] = and i64 [[TMP48]], -3
+// CHECK-NEXT: br label [[OMP_TYPE_END22]]
+// CHECK: omp.type.to.else20:
+// CHECK-NEXT: [[TMP54:%.*]] = icmp eq i64 [[TMP49]], 2
+// CHECK-NEXT: br i1 [[TMP54]], label [[OMP_TYPE_FROM21:%.*]], label [[OMP_TYPE_END22]]
+// CHECK: omp.type.from21:
+// CHECK-NEXT: [[TMP55:%.*]] = and i64 [[TMP48]], -2
+// CHECK-NEXT: br label [[OMP_TYPE_END22]]
+// CHECK: omp.type.end22:
+// CHECK-NEXT: [[OMP_MAPTYPE23:%.*]] = phi i64 [ [[TMP51]], [[OMP_TYPE_ALLOC17]] ], [ [[TMP53]], [[OMP_TYPE_TO19]] ], [ [[TMP55]], [[OMP_TYPE_FROM21]] ], [ [[TMP48]], [[OMP_TYPE_TO_ELSE20]] ]
+// CHECK-NEXT: [[TMP56:%.*]] = and i64 [[TMP4]], 5132
+// CHECK-NEXT: [[OMP_MAPTYPE_WITH_MODIFIERS24:%.*]] = or i64 [[OMP_MAPTYPE23]], [[TMP56]]
+// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[H]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS24]], ptr null)
// CHECK-NEXT: [[OMP_ARRAYMAP_NEXT]] = getelementptr [[STRUCT_D]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 1
// CHECK-NEXT: [[OMP_ARRAYMAP_ISDONE:%.*]] = icmp eq ptr [[OMP_ARRAYMAP_NEXT]], [[TMP7]]
// CHECK-NEXT: br i1 [[OMP_ARRAYMAP_ISDONE]], label [[OMP_ARRAYMAP_EXIT:%.*]], label [[OMP_ARRAYMAP_BODY]]
// CHECK: omp.arraymap.exit:
-// CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY22:%.*]] = icmp sgt i64 [[TMP6]], 1
-// CHECK-NEXT: [[TMP53:%.*]] = and i64 [[TMP4]], 8
-// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP53]], 0
-// CHECK-NEXT: [[TMP54:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY22]], [[DOTOMP_ARRAY__DEL__DELETE]]
-// CHECK-NEXT: br i1 [[TMP54]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]]
+// CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY25:%.*]] = icmp sgt i64 [[TMP6]], 1
+// CHECK-NEXT: [[TMP57:%.*]] = and i64 [[TMP4]], 8
+// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP57]], 0
+// CHECK-NEXT: [[TMP58:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY25]], [[DOTOMP_ARRAY__DEL__DELETE]]
+// CHECK-NEXT: br i1 [[TMP58]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]]
// CHECK: .omp.array..del:
-// CHECK-NEXT: [[TMP55:%.*]] = mul nuw i64 [[TMP6]], 12
-// CHECK-NEXT: [[TMP56:%.*]] = and i64 [[TMP4]], -4
-// CHECK-NEXT: [[TMP57:%.*]] = or i64 [[TMP56]], 512
-// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP55]], i64 [[TMP57]], ptr [[TMP5]])
+// CHECK-NEXT: [[TMP59:%.*]] = mul nuw i64 [[TMP6]], 12
+// CHECK-NEXT: [[TMP60:%.*]] = and i64 [[TMP4]], -4
+// CHECK-NEXT: [[TMP61:%.*]] = or i64 [[TMP60]], 512
+// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP59]], i64 [[TMP61]], ptr [[TMP5]])
// CHECK-NEXT: br label [[OMP_DONE]]
// CHECK: omp.done:
// CHECK-NEXT: ret void
@@ -303,21 +311,23 @@ void foo() {
// CHECK-NEXT: br label [[OMP_TYPE_END]]
// CHECK: omp.type.end:
// CHECK-NEXT: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP20]], [[OMP_TYPE_ALLOC]] ], [ [[TMP22]], [[OMP_TYPE_TO]] ], [ [[TMP24]], [[OMP_TYPE_FROM]] ], [ [[TMP17]], [[OMP_TYPE_TO_ELSE]] ]
-// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[A]], i64 4, i64 [[OMP_MAPTYPE]], ptr null)
+// CHECK-NEXT: [[TMP25:%.*]] = and i64 [[TMP4]], 5132
+// CHECK-NEXT: [[OMP_MAPTYPE_WITH_MODIFIERS:%.*]] = or i64 [[OMP_MAPTYPE]], [[TMP25]]
+// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[A]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS]], ptr null)
// CHECK-NEXT: [[OMP_ARRAYMAP_NEXT]] = getelementptr [[STRUCT_C]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 1
// CHECK-NEXT: [[OMP_ARRAYMAP_ISDONE:%.*]] = icmp eq ptr [[OMP_ARRAYMAP_NEXT]], [[TMP7]]
// CHECK-NEXT: br i1 [[OMP_ARRAYMAP_ISDONE]], label [[OMP_ARRAYMAP_EXIT:%.*]], label [[OMP_ARRAYMAP_BODY]]
// CHECK: omp.arraymap.exit:
// CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY1:%.*]] = icmp sgt i64 [[TMP6]], 1
-// CHECK-NEXT: [[TMP25:%.*]] = and i64 [[TMP4]], 8
-// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP25]], 0
-// CHECK-NEXT: [[TMP26:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY1]], [[DOTOMP_ARRAY__DEL__DELETE]]
-// CHECK-NEXT: br i1 [[TMP26]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]]
+// CHECK-NEXT: [[TMP26:%.*]] = and i64 [[TMP4]], 8
+// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP26]], 0
+// CHECK-NEXT: [[TMP27:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY1]], [[DOTOMP_ARRAY__DEL__DELETE]]
+// CHECK-NEXT: br i1 [[TMP27]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]]
// CHECK: .omp.array..del:
-// CHECK-NEXT: [[TMP27:%.*]] = mul nuw i64 [[TMP6]], 4
-// CHECK-NEXT: [[TMP28:%.*]] = and i64 [[TMP4]], -4
-// CHECK-NEXT: [[TMP29:%.*]] = or i64 [[TMP28]], 512
-// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP27]], i64 [[TMP29]], ptr [[TMP5]])
+// CHECK-NEXT: [[TMP28:%.*]] = mul nuw i64 [[TMP6]], 4
+// CHECK-NEXT: [[TMP29:%.*]] = and i64 [[TMP4]], -4
+// CHECK-NEXT: [[TMP30:%.*]] = or i64 [[TMP29]], 512
+// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP28]], i64 [[TMP30]], ptr [[TMP5]])
// CHECK-NEXT: br label [[OMP_DONE]]
// CHECK: omp.done:
// CHECK-NEXT: ret void
diff --git a/clang/test/OpenMP/target_map_array_section_of_structs_with_nested_mapper_codegen.cpp b/clang/test/OpenMP/target_map_array_section_of_structs_with_nested_mapper_codegen.cpp
index daee73d80480e..f72a2e68c0d3e 100644
--- a/clang/test/OpenMP/target_map_array_section_of_structs_with_nested_mapper_codegen.cpp
+++ b/clang/test/OpenMP/target_map_array_section_of_structs_with_nested_mapper_codegen.cpp
@@ -136,7 +136,7 @@ void foo() {
// CHECK-NEXT: [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP2]], [[TMP7]]
// CHECK-NEXT: br i1 [[OMP_ARRAYMAP_ISEMPTY]], label [[OMP_DONE:%.*]], label [[OMP_ARRAYMAP_BODY:%.*]]
// CHECK: omp.arraymap.body:
-// CHECK-NEXT: [[OMP_ARRAYMAP_PTRCURRENT:%.*]] = phi ptr [ [[TMP2]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END20:%.*]] ]
+// CHECK-NEXT: [[OMP_ARRAYMAP_PTRCURRENT:%.*]] = phi ptr [ [[TMP2]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END22:%.*]] ]
// CHECK-NEXT: [[E:%.*]] = getelementptr inbounds nuw [[STRUCT_D]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
// CHECK-NEXT: [[F:%.*]] = getelementptr inbounds nuw [[STRUCT_D]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 1
// CHECK-NEXT: [[H:%.*]] = getelementptr inbounds nuw [[STRUCT_D]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 2
@@ -167,87 +167,95 @@ void foo() {
// CHECK-NEXT: br label [[OMP_TYPE_END]]
// CHECK: omp.type.end:
// CHECK-NEXT: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP24]], [[OMP_TYPE_ALLOC]] ], [ [[TMP26]], [[OMP_TYPE_TO]] ], [ [[TMP28]], [[OMP_TYPE_FROM]] ], [ [[TMP21]], [[OMP_TYPE_TO_ELSE]] ]
-// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 [[TMP18]], i64 [[OMP_MAPTYPE]], ptr null)
-// CHECK-NEXT: [[TMP29:%.*]] = add nuw i64 281474976711171, [[TMP20]]
-// CHECK-NEXT: [[TMP30:%.*]] = and i64 [[TMP4]], 3
-// CHECK-NEXT: [[TMP31:%.*]] = icmp eq i64 [[TMP30]], 0
-// CHECK-NEXT: br i1 [[TMP31]], label [[OMP_TYPE_ALLOC1:%.*]], label [[OMP_TYPE_ALLOC_ELSE2:%.*]]
+// CHECK-NEXT: [[TMP29:%.*]] = and i64 [[TMP4]], 5132
+// CHECK-NEXT: [[OMP_MAPTYPE_WITH_MODIFIERS:%.*]] = or i64 [[OMP_MAPTYPE]], [[TMP29]]
+// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 [[TMP18]], i64 [[OMP_MAPTYPE_WITH_MODIFIERS]], ptr null)
+// CHECK-NEXT: [[TMP30:%.*]] = add nuw i64 281474976711171, [[TMP20]]
+// CHECK-NEXT: [[TMP31:%.*]] = and i64 [[TMP4]], 3
+// CHECK-NEXT: [[TMP32:%.*]] = icmp eq i64 [[TMP31]], 0
+// CHECK-NEXT: br i1 [[TMP32]], label [[OMP_TYPE_ALLOC1:%.*]], label [[OMP_TYPE_ALLOC_ELSE2:%.*]]
// CHECK: omp.type.alloc1:
-// CHECK-NEXT: [[TMP32:%.*]] = and i64 [[TMP29]], -4
+// CHECK-NEXT: [[TMP33:%.*]] = and i64 [[TMP30]], -4
// CHECK-NEXT: br label [[OMP_TYPE_END6:%.*]]
// CHECK: omp.type.alloc.else2:
-// CHECK-NEXT: [[TMP33:%.*]] = icmp eq i64 [[TMP30]], 1
-// CHECK-NEXT: br i1 [[TMP33]], label [[OMP_TYPE_TO3:%.*]], label [[OMP_TYPE_TO_ELSE4:%.*]]
+// CHECK-NEXT: [[TMP34:%.*]] = icmp eq i64 [[TMP31]], 1
+// CHECK-NEXT: br i1 [[TMP34]], label [[OMP_TYPE_TO3:%.*]], label [[OMP_TYPE_TO_ELSE4:%.*]]
// CHECK: omp.type.to3:
-// CHECK-NEXT: [[TMP34:%.*]] = and i64 [[TMP29]], -3
+// CHECK-NEXT: [[TMP35:%.*]] = and i64 [[TMP30]], -3
// CHECK-NEXT: br label [[OMP_TYPE_END6]]
// CHECK: omp.type.to.else4:
-// CHECK-NEXT: [[TMP35:%.*]] = icmp eq i64 [[TMP30]], 2
-// CHECK-NEXT: br i1 [[TMP35]], label [[OMP_TYPE_FROM5:%.*]], label [[OMP_TYPE_END6]]
+// CHECK-NEXT: [[TMP36:%.*]] = icmp eq i64 [[TMP31]], 2
+// CHECK-NEXT: br i1 [[TMP36]], label [[OMP_TYPE_FROM5:%.*]], label [[OMP_TYPE_END6]]
// CHECK: omp.type.from5:
-// CHECK-NEXT: [[TMP36:%.*]] = and i64 [[TMP29]], -2
+// CHECK-NEXT: [[TMP37:%.*]] = and i64 [[TMP30]], -2
// CHECK-NEXT: br label [[OMP_TYPE_END6]]
// CHECK: omp.type.end6:
-// CHECK-NEXT: [[OMP_MAPTYPE7:%.*]] = phi i64 [ [[TMP32]], [[OMP_TYPE_ALLOC1]] ], [ [[TMP34]], [[OMP_TYPE_TO3]] ], [ [[TMP36]], [[OMP_TYPE_FROM5]] ], [ [[TMP29]], [[OMP_TYPE_TO_ELSE4]] ]
-// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 4, i64 [[OMP_MAPTYPE7]], ptr null)
-// CHECK-NEXT: [[TMP37:%.*]] = add nuw i64 281474976711171, [[TMP20]]
-// CHECK-NEXT: [[TMP38:%.*]] = and i64 [[TMP4]], 3
-// CHECK-NEXT: [[TMP39:%.*]] = icmp eq i64 [[TMP38]], 0
-// CHECK-NEXT: br i1 [[TMP39]], label [[OMP_TYPE_ALLOC8:%.*]], label [[OMP_TYPE_ALLOC_ELSE9:%.*]]
-// CHECK: omp.type.alloc8:
-// CHECK-NEXT: [[TMP40:%.*]] = and i64 [[TMP37]], -4
-// CHECK-NEXT: br label [[OMP_TYPE_END13:%.*]]
-// CHECK: omp.type.alloc.else9:
-// CHECK-NEXT: [[TMP41:%.*]] = icmp eq i64 [[TMP38]], 1
-// CHECK-NEXT: br i1 [[TMP41]], label [[OMP_TYPE_TO10:%.*]], label [[OMP_TYPE_TO_ELSE11:%.*]]
-// CHECK: omp.type.to10:
-// CHECK-NEXT: [[TMP42:%.*]] = and i64 [[TMP37]], -3
-// CHECK-NEXT: br label [[OMP_TYPE_END13]]
-// CHECK: omp.type.to.else11:
-// CHECK-NEXT: [[TMP43:%.*]] = icmp eq i64 [[TMP38]], 2
-// CHECK-NEXT: br i1 [[TMP43]], label [[OMP_TYPE_FROM12:%.*]], label [[OMP_TYPE_END13]]
-// CHECK: omp.type.from12:
-// CHECK-NEXT: [[TMP44:%.*]] = and i64 [[TMP37]], -2
-// CHECK-NEXT: br label [[OMP_TYPE_END13]]
-// CHECK: omp.type.end13:
-// CHECK-NEXT: [[OMP_MAPTYPE14:%.*]] = phi i64 [ [[TMP40]], [[OMP_TYPE_ALLOC8]] ], [ [[TMP42]], [[OMP_TYPE_TO10]] ], [ [[TMP44]], [[OMP_TYPE_FROM12]] ], [ [[TMP37]], [[OMP_TYPE_TO_ELSE11]] ]
-// CHECK-NEXT: call void @.omp_mapper._ZTS1C.default(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[F]], i64 4, i64 [[OMP_MAPTYPE14]], ptr null) #[[ATTR3]]
-// CHECK-NEXT: [[TMP45:%.*]] = add nuw i64 281474976711171, [[TMP20]]
-// CHECK-NEXT: [[TMP46:%.*]] = and i64 [[TMP4]], 3
-// CHECK-NEXT: [[TMP47:%.*]] = icmp eq i64 [[TMP46]], 0
-// CHECK-NEXT: br i1 [[TMP47]], label [[OMP_TYPE_ALLOC15:%.*]], label [[OMP_TYPE_ALLOC_ELSE16:%.*]]
-// CHECK: omp.type.alloc15:
-// CHECK-NEXT: [[TMP48:%.*]] = and i64 [[TMP45]], -4
-// CHECK-NEXT: br label [[OMP_TYPE_END20]]
-// CHECK: omp.type.alloc.else16:
-// CHECK-NEXT: [[TMP49:%.*]] = icmp eq i64 [[TMP46]], 1
-// CHECK-NEXT: br i1 [[TMP49]], label [[OMP_TYPE_TO17:%.*]], label [[OMP_TYPE_TO_ELSE18:%.*]]
-// CHECK: omp.type.to17:
-// CHECK-NEXT: [[TMP50:%.*]] = and i64 [[TMP45]], -3
-// CHECK-NEXT: br label [[OMP_TYPE_END20]]
-// CHECK: omp.type.to.else18:
-// CHECK-NEXT: [[TMP51:%.*]] = icmp eq i64 [[TMP46]], 2
-// CHECK-NEXT: br i1 [[TMP51]], label [[OMP_TYPE_FROM19:%.*]], label [[OMP_TYPE_END20]]
-// CHECK: omp.type.from19:
-// CHECK-NEXT: [[TMP52:%.*]] = and i64 [[TMP45]], -2
-// CHECK-NEXT: br label [[OMP_TYPE_END20]]
-// CHECK: omp.type.end20:
-// CHECK-NEXT: [[OMP_MAPTYPE21:%.*]] = phi i64 [ [[TMP48]], [[OMP_TYPE_ALLOC15]] ], [ [[TMP50]], [[OMP_TYPE_TO17]] ], [ [[TMP52]], [[OMP_TYPE_FROM19]] ], [ [[TMP45]], [[OMP_TYPE_TO_ELSE18]] ]
-// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[H]], i64 4, i64 [[OMP_MAPTYPE21]], ptr null)
+// CHECK-NEXT: [[OMP_MAPTYPE7:%.*]] = phi i64 [ [[TMP33]], [[OMP_TYPE_ALLOC1]] ], [ [[TMP35]], [[OMP_TYPE_TO3]] ], [ [[TMP37]], [[OMP_TYPE_FROM5]] ], [ [[TMP30]], [[OMP_TYPE_TO_ELSE4]] ]
+// CHECK-NEXT: [[TMP38:%.*]] = and i64 [[TMP4]], 5132
+// CHECK-NEXT: [[OMP_MAPTYPE_WITH_MODIFIERS8:%.*]] = or i64 [[OMP_MAPTYPE7]], [[TMP38]]
+// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS8]], ptr null)
+// CHECK-NEXT: [[TMP39:%.*]] = add nuw i64 281474976711171, [[TMP20]]
+// CHECK-NEXT: [[TMP40:%.*]] = and i64 [[TMP4]], 3
+// CHECK-NEXT: [[TMP41:%.*]] = icmp eq i64 [[TMP40]], 0
+// CHECK-NEXT: br i1 [[TMP41]], label [[OMP_TYPE_ALLOC9:%.*]], label [[OMP_TYPE_ALLOC_ELSE10:%.*]]
+// CHECK: omp.type.alloc9:
+// CHECK-NEXT: [[TMP42:%.*]] = and i64 [[TMP39]], -4
+// CHECK-NEXT: br label [[OMP_TYPE_END14:%.*]]
+// CHECK: omp.type.alloc.else10:
+// CHECK-NEXT: [[TMP43:%.*]] = icmp eq i64 [[TMP40]], 1
+// CHECK-NEXT: br i1 [[TMP43]], label [[OMP_TYPE_TO11:%.*]], label [[OMP_TYPE_TO_ELSE12:%.*]]
+// CHECK: omp.type.to11:
+// CHECK-NEXT: [[TMP44:%.*]] = and i64 [[TMP39]], -3
+// CHECK-NEXT: br label [[OMP_TYPE_END14]]
+// CHECK: omp.type.to.else12:
+// CHECK-NEXT: [[TMP45:%.*]] = icmp eq i64 [[TMP40]], 2
+// CHECK-NEXT: br i1 [[TMP45]], label [[OMP_TYPE_FROM13:%.*]], label [[OMP_TYPE_END14]]
+// CHECK: omp.type.from13:
+// CHECK-NEXT: [[TMP46:%.*]] = and i64 [[TMP39]], -2
+// CHECK-NEXT: br label [[OMP_TYPE_END14]]
+// CHECK: omp.type.end14:
+// CHECK-NEXT: [[OMP_MAPTYPE15:%.*]] = phi i64 [ [[TMP42]], [[OMP_TYPE_ALLOC9]] ], [ [[TMP44]], [[OMP_TYPE_TO11]] ], [ [[TMP46]], [[OMP_TYPE_FROM13]] ], [ [[TMP39]], [[OMP_TYPE_TO_ELSE12]] ]
+// CHECK-NEXT: [[TMP47:%.*]] = and i64 [[TMP4]], 5132
+// CHECK-NEXT: [[OMP_MAPTYPE_WITH_MODIFIERS16:%.*]] = or i64 [[OMP_MAPTYPE15]], [[TMP47]]
+// CHECK-NEXT: call void @.omp_mapper._ZTS1C.default(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[F]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS16]], ptr null) #[[ATTR3]]
+// CHECK-NEXT: [[TMP48:%.*]] = add nuw i64 281474976711171, [[TMP20]]
+// CHECK-NEXT: [[TMP49:%.*]] = and i64 [[TMP4]], 3
+// CHECK-NEXT: [[TMP50:%.*]] = icmp eq i64 [[TMP49]], 0
+// CHECK-NEXT: br i1 [[TMP50]], label [[OMP_TYPE_ALLOC17:%.*]], label [[OMP_TYPE_ALLOC_ELSE18:%.*]]
+// CHECK: omp.type.alloc17:
+// CHECK-NEXT: [[TMP51:%.*]] = and i64 [[TMP48]], -4
+// CHECK-NEXT: br label [[OMP_TYPE_END22]]
+// CHECK: omp.type.alloc.else18:
+// CHECK-NEXT: [[TMP52:%.*]] = icmp eq i64 [[TMP49]], 1
+// CHECK-NEXT: br i1 [[TMP52]], label [[OMP_TYPE_TO19:%.*]], label [[OMP_TYPE_TO_ELSE20:%.*]]
+// CHECK: omp.type.to19:
+// CHECK-NEXT: [[TMP53:%.*]] = and i64 [[TMP48]], -3
+// CHECK-NEXT: br label [[OMP_TYPE_END22]]
+// CHECK: omp.type.to.else20:
+// CHECK-NEXT: [[TMP54:%.*]] = icmp eq i64 [[TMP49]], 2
+// CHECK-NEXT: br i1 [[TMP54]], label [[OMP_TYPE_FROM21:%.*]], label [[OMP_TYPE_END22]]
+// CHECK: omp.type.from21:
+// CHECK-NEXT: [[TMP55:%.*]] = and i64 [[TMP48]], -2
+// CHECK-NEXT: br label [[OMP_TYPE_END22]]
+// CHECK: omp.type.end22:
+// CHECK-NEXT: [[OMP_MAPTYPE23:%.*]] = phi i64 [ [[TMP51]], [[OMP_TYPE_ALLOC17]] ], [ [[TMP53]], [[OMP_TYPE_TO19]] ], [ [[TMP55]], [[OMP_TYPE_FROM21]] ], [ [[TMP48]], [[OMP_TYPE_TO_ELSE20]] ]
+// CHECK-NEXT: [[TMP56:%.*]] = and i64 [[TMP4]], 5132
+// CHECK-NEXT: [[OMP_MAPTYPE_WITH_MODIFIERS24:%.*]] = or i64 [[OMP_MAPTYPE23]], [[TMP56]]
+// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[H]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS24]], ptr null)
// CHECK-NEXT: [[OMP_ARRAYMAP_NEXT]] = getelementptr [[STRUCT_D]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 1
// CHECK-NEXT: [[OMP_ARRAYMAP_ISDONE:%.*]] = icmp eq ptr [[OMP_ARRAYMAP_NEXT]], [[TMP7]]
// CHECK-NEXT: br i1 [[OMP_ARRAYMAP_ISDONE]], label [[OMP_ARRAYMAP_EXIT:%.*]], label [[OMP_ARRAYMAP_BODY]]
// CHECK: omp.arraymap.exit:
-// CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY22:%.*]] = icmp sgt i64 [[TMP6]], 1
-// CHECK-NEXT: [[TMP53:%.*]] = and i64 [[TMP4]], 8
-// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP53]], 0
-// CHECK-NEXT: [[TMP54:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY22]], [[DOTOMP_ARRAY__DEL__DELETE]]
-// CHECK-NEXT: br i1 [[TMP54]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]]
+// CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY25:%.*]] = icmp sgt i64 [[TMP6]], 1
+// CHECK-NEXT: [[TMP57:%.*]] = and i64 [[TMP4]], 8
+// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP57]], 0
+// CHECK-NEXT: [[TMP58:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY25]], [[DOTOMP_ARRAY__DEL__DELETE]]
+// CHECK-NEXT: br i1 [[TMP58]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]]
// CHECK: .omp.array..del:
-// CHECK-NEXT: [[TMP55:%.*]] = mul nuw i64 [[TMP6]], 12
-// CHECK-NEXT: [[TMP56:%.*]] = and i64 [[TMP4]], -4
-// CHECK-NEXT: [[TMP57:%.*]] = or i64 [[TMP56]], 512
-// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP55]], i64 [[TMP57]], ptr [[TMP5]])
+// CHECK-NEXT: [[TMP59:%.*]] = mul nuw i64 [[TMP6]], 12
+// CHECK-NEXT: [[TMP60:%.*]] = and i64 [[TMP4]], -4
+// CHECK-NEXT: [[TMP61:%.*]] = or i64 [[TMP60]], 512
+// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP59]], i64 [[TMP61]], ptr [[TMP5]])
// CHECK-NEXT: br label [[OMP_DONE]]
// CHECK: omp.done:
// CHECK-NEXT: ret void
@@ -300,21 +308,23 @@ void foo() {
// CHECK-NEXT: br label [[OMP_TYPE_END]]
// CHECK: omp.type.end:
// CHECK-NEXT: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP20]], [[OMP_TYPE_ALLOC]] ], [ [[TMP22]], [[OMP_TYPE_TO]] ], [ [[TMP24]], [[OMP_TYPE_FROM]] ], [ [[TMP17]], [[OMP_TYPE_TO_ELSE]] ]
-// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[A]], i64 4, i64 [[OMP_MAPTYPE]], ptr null)
+// CHECK-NEXT: [[TMP25:%.*]] = and i64 [[TMP4]], 5132
+// CHECK-NEXT: [[OMP_MAPTYPE_WITH_MODIFIERS:%.*]] = or i64 [[OMP_MAPTYPE]], [[TMP25]]
+// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[A]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS]], ptr null)
// CHECK-NEXT: [[OMP_ARRAYMAP_NEXT]] = getelementptr [[STRUCT_C]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 1
// CHECK-NEXT: [[OMP_ARRAYMAP_ISDONE:%.*]] = icmp eq ptr [[OMP_ARRAYMAP_NEXT]], [[TMP7]]
// CHECK-NEXT: br i1 [[OMP_ARRAYMAP_ISDONE]], label [[OMP_ARRAYMAP_EXIT:%.*]], label [[OMP_ARRAYMAP_BODY]]
// CHECK: omp.arraymap.exit:
// CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY1:%.*]] = icmp sgt i64 [[TMP6]], 1
-// CHECK-NEXT: [[TMP25:%.*]] = and i64 [[TMP4]], 8
-// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP25]], 0
-// CHECK-NEXT: [[TMP26:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY1]], [[DOTOMP_ARRAY__DEL__DELETE]]
-// CHECK-NEXT: br i1 [[TMP26]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]]
+// CHECK-NEXT: [[TMP26:%.*]] = and i64 [[TMP4]], 8
+// CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP26]], 0
+// CHECK-NEXT: [[TMP27:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY1]], [[DOTOMP_ARRAY__DEL__DELETE]]
+// CHECK-NEXT: br i1 [[TMP27]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]]
// CHECK: .omp.array..del:
-// CHECK-NEXT: [[TMP27:%.*]] = mul nuw i64 [[TMP6]], 4
-// CHECK-NEXT: [[TMP28:%.*]] = and i64 [[TMP4]], -4
-// CHECK-NEXT: [[TMP29:%.*]] = or i64 [[TMP28]], 512
-// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP27]], i64 [[TMP29]], ptr [[TMP5]])
+// CHECK-NEXT: [[TMP28:%.*]] = mul nuw i64 [[TMP6]], 4
+// CHECK-NEXT: [[TMP29:%.*]] = and i64 [[TMP4]], -4
+// CHECK-NEXT: [[TMP30:%.*]] = or i64 [[TMP29]], 512
+// CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP28]], i64 [[TMP30]], ptr [[TMP5]])
// CHECK-NEXT: br label [[OMP_DONE]]
// CHECK: omp.done:
// CHECK-NEXT: ret void
diff --git a/clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp b/clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp
index d90d4fa3fcfbb..82509ace56eec 100644
--- a/clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp
+++ b/clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp
@@ -39,8 +39,8 @@ void foo(S2 *arr) {
#endif
//.
-// CHECK: @.offload_sizes = private unnamed_addr constant [1 x i64] [i64 32]
-// CHECK: @.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 [[#0x1]]]
+// CHECK: @.offload_sizes = private unnamed_addr constant [2 x i64] [i64 32, i64 8]
+// CHECK: @.offload_maptypes = private unnamed_addr constant [2 x i64] [i64 [[#0x1]], i64 [[#0x4000]]]
//.
// CHECK-LABEL: define {{[^@]+}}@_Z3fooP2S2
// CHECK-SAME: (ptr noundef [[ARR:%.*]]) #[[ATTR0:[0-9]+]] {
@@ -48,175 +48,191 @@ void foo(S2 *arr) {
// CHECK: store ptr [[ARR]], ptr [[ARR_ADDR:%.*]], align 8
// CHECK: [[TMP0:%.*]] = load ptr, ptr [[ARR_ADDR]], align 8
// CHECK: [[TMP1:%.*]] = load ptr, ptr [[ARR_ADDR]], align 8
-// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds [[STRUCT_S2:%.*]], ptr [[TMP1]], i64 0
-// CHECK: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0
+// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds nuw [[STRUCT_S2:%.*]], ptr [[TMP1]], i64 0
+// CHECK: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0
// CHECK: store ptr [[TMP0]], ptr [[TMP2]], align 8
-// CHECK: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0
+// CHECK: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0
// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP3]], align 8
-// CHECK: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0
+// CHECK: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0
// CHECK: store ptr @.omp_mapper._ZTS2S2.default, ptr [[TMP4]], align 8
-// CHECK: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
-// CHECK: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
-// CHECK: call void @__tgt_target_data_begin_mapper(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 1, ptr [[TMP5]], ptr [[TMP6]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr [[DOTOFFLOAD_MAPPERS]])
+// CHECK: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1
+// CHECK: store ptr [[ARR_ADDR]], ptr [[TMP5]], align 8
+// CHECK: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 1
+// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP6]], align 8
+// CHECK: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 1
+// CHECK: store ptr null, ptr [[TMP7]], align 8
+// CHECK: [[TMP8:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// CHECK: [[TMP9:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// CHECK: call void @__tgt_target_data_begin_mapper(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 2, ptr [[TMP8]], ptr [[TMP9]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr [[DOTOFFLOAD_MAPPERS]])
// CHECK: ret void
//
//
// CHECK-LABEL: define {{[^@]+}}@.omp_mapper._ZTS2S2.default
// CHECK-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]], ptr noundef [[TMP2:%.*]], i64 noundef [[TMP3:%.*]], i64 noundef [[TMP4:%.*]], ptr noundef [[TMP5:%.*]]) #[[ATTR1:[0-9]+]] {
// CHECK: entry:
-// CHECK: store ptr [[TMP0]], ptr [[DOTADDR:%.*]], align 8
-// CHECK: store ptr [[TMP1]], ptr [[DOTADDR1:%.*]], align 8
-// CHECK: store ptr [[TMP2]], ptr [[DOTADDR2:%.*]], align 8
-// CHECK: store i64 [[TMP3]], ptr [[DOTADDR3:%.*]], align 8
-// CHECK: store i64 [[TMP4]], ptr [[DOTADDR4:%.*]], align 8
-// CHECK: store ptr [[TMP5]], ptr [[DOTADDR5:%.*]], align 8
-// CHECK: [[TMP6:%.*]] = load i64, ptr [[DOTADDR3]], align 8
-// CHECK: [[TMP7:%.*]] = load ptr, ptr [[DOTADDR]], align 8
-// CHECK: [[TMP8:%.*]] = load ptr, ptr [[DOTADDR1]], align 8
-// CHECK: [[TMP9:%.*]] = load ptr, ptr [[DOTADDR2]], align 8
-// CHECK: [[TMP10:%.*]] = udiv exact i64 [[TMP6]], 16
-// CHECK: [[TMP11:%.*]] = getelementptr [[STRUCT_S2:%.*]], ptr [[TMP9]], i64 [[TMP10]]
-// CHECK: [[TMP12:%.*]] = load i64, ptr [[DOTADDR4]], align 8
-// CHECK: [[TMP13:%.*]] = load ptr, ptr [[DOTADDR5]], align 8
-// CHECK: [[OMP_ARRAYINIT_ISARRAY:%.*]] = icmp sgt i64 [[TMP10]], 1
-// CHECK: [[TMP14:%.*]] = and i64 [[TMP12]], 8
-// CHECK: [[TMP15:%.*]] = icmp ne ptr [[TMP8]], [[TMP9]]
-// CHECK: [[TMP16:%.*]] = and i64 [[TMP12]], 16
-// CHECK: [[TMP17:%.*]] = icmp ne i64 [[TMP16]], 0
-// CHECK: [[TMP18:%.*]] = and i1 [[TMP15]], [[TMP17]]
-// CHECK: [[TMP19:%.*]] = or i1 [[OMP_ARRAYINIT_ISARRAY]], [[TMP18]]
-// CHECK: [[DOTOMP_ARRAY__INIT__DELETE:%.*]] = icmp eq i64 [[TMP14]], 0
-// CHECK: [[TMP20:%.*]] = and i1 [[TMP19]], [[DOTOMP_ARRAY__INIT__DELETE]]
-// CHECK: br i1 [[TMP20]], label [[DOTOMP_ARRAY__INIT:%.*]], label [[OMP_ARRAYMAP_HEAD:%.*]]
+// CHECK: [[TMP6:%.*]] = udiv exact i64 [[TMP3]], 16
+// CHECK: [[TMP7:%.*]] = getelementptr [[STRUCT_S2:%.*]], ptr [[TMP2]], i64 [[TMP6]]
+// CHECK: [[OMP_ARRAYINIT_ISARRAY:%.*]] = icmp sgt i64 [[TMP6]], 1
+// CHECK: [[TMP8:%.*]] = and i64 [[TMP4]], 8
+// CHECK: [[TMP9:%.*]] = icmp ne ptr [[TMP1]], [[TMP2]]
+// CHECK: [[TMP10:%.*]] = or i1 [[OMP_ARRAYINIT_ISARRAY]], [[TMP9]]
+// CHECK: [[DOTOMP_ARRAY__INIT__DELETE:%.*]] = icmp eq i64 [[TMP8]], 0
+// CHECK: [[TMP11:%.*]] = and i1 [[TMP10]], [[DOTOMP_ARRAY__INIT__DELETE]]
+// CHECK: br i1 [[TMP11]], label [[DOTOMP_ARRAY__INIT:%.*]], label [[OMP_ARRAYMAP_HEAD:%.*]]
// CHECK: .omp.array..init:
-// CHECK: [[TMP21:%.*]] = mul nuw i64 [[TMP10]], 16
-// CHECK: [[TMP22:%.*]] = and i64 [[TMP12]], -4
-// CHECK: [[TMP23:%.*]] = or i64 [[TMP22]], 512
-// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP7]], ptr [[TMP8]], ptr [[TMP9]], i64 [[TMP21]], i64 [[TMP23]], ptr [[TMP13]])
+// CHECK: [[TMP12:%.*]] = mul nuw i64 [[TMP6]], 16
+// CHECK: [[TMP13:%.*]] = and i64 [[TMP4]], -4
+// CHECK: [[TMP14:%.*]] = or i64 [[TMP13]], 512
+// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP12]], i64 [[TMP14]], ptr [[TMP5]])
// CHECK: br label [[OMP_ARRAYMAP_HEAD]]
// CHECK: omp.arraymap.head:
-// CHECK: [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP9]], [[TMP11]]
+// CHECK: [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP2]], [[TMP7]]
// CHECK: br i1 [[OMP_ARRAYMAP_ISEMPTY]], label [[OMP_DONE:%.*]], label [[OMP_ARRAYMAP_BODY:%.*]]
// CHECK: omp.arraymap.body:
-// CHECK: [[OMP_ARRAYMAP_PTRCURRENT:%.*]] = phi ptr [ [[TMP9]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END28:%.*]] ]
-// CHECK: [[Z:%.*]] = getelementptr inbounds [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 1
-// CHECK: [[S1P:%.*]] = getelementptr inbounds [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
-// CHECK: [[S1P6:%.*]] = getelementptr inbounds [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
-// CHECK: [[TMP24:%.*]] = load ptr, ptr [[S1P6]], align 8
-// CHECK: [[X:%.*]] = getelementptr inbounds [[STRUCT_S1:%.*]], ptr [[TMP24]], i32 0, i32 0
-// CHECK: [[S1P7:%.*]] = getelementptr inbounds [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
-// CHECK: [[S1P8:%.*]] = getelementptr inbounds [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
-// CHECK: [[TMP25:%.*]] = load ptr, ptr [[S1P8]], align 8
-// CHECK: [[Y:%.*]] = getelementptr inbounds [[STRUCT_S1]], ptr [[TMP25]], i32 0, i32 1
-// CHECK: [[TMP26:%.*]] = getelementptr i32, ptr [[Z]], i32 1
-// CHECK: [[TMP27:%.*]] = ptrtoint ptr [[TMP26]] to i64
-// CHECK: [[TMP28:%.*]] = ptrtoint ptr [[S1P]] to i64
-// CHECK: [[TMP29:%.*]] = sub i64 [[TMP27]], [[TMP28]]
-// CHECK: [[TMP30:%.*]] = sdiv exact i64 [[TMP29]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)
-// CHECK: [[TMP31:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP7]])
-// CHECK: [[TMP32:%.*]] = shl i64 [[TMP31]], 48
-// CHECK: [[TMP33:%.*]] = add nuw i64 0, [[TMP32]]
-// CHECK: [[TMP34:%.*]] = and i64 [[TMP12]], 3
-// CHECK: [[TMP35:%.*]] = icmp eq i64 [[TMP34]], 0
-// CHECK: br i1 [[TMP35]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]]
+// CHECK: [[OMP_ARRAYMAP_PTRCURRENT:%.*]] = phi ptr [ [[TMP2]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END33:%.*]] ]
+// CHECK: [[Z:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 1
+// CHECK: [[S1P:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
+// CHECK: [[TMP15:%.*]] = load ptr, ptr [[S1P]], align 8
+// CHECK: [[S1P1:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
+// CHECK: [[TMP16:%.*]] = load ptr, ptr [[S1P1]], align 8
+// CHECK: [[X:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP16]], i32 0, i32 0
+// CHECK: [[S1P2:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
+// CHECK: [[TMP17:%.*]] = load ptr, ptr [[S1P2]], align 8
+// CHECK: [[S1P3:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
+// CHECK: [[TMP18:%.*]] = load ptr, ptr [[S1P3]], align 8
+// CHECK: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S1]], ptr [[TMP18]], i32 0, i32 1
+// CHECK: [[TMP19:%.*]] = getelementptr i32, ptr [[Y]], i32 1
+// CHECK: [[TMP20:%.*]] = ptrtoaddr ptr [[TMP19]] to i64
+// CHECK: [[TMP21:%.*]] = ptrtoaddr ptr [[X]] to i64
+// CHECK: [[TMP22:%.*]] = sub i64 [[TMP20]], [[TMP21]]
+// CHECK: [[TMP23:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP0]])
+// CHECK: [[TMP24:%.*]] = shl i64 [[TMP23]], 48
+// CHECK: [[TMP25:%.*]] = add nuw i64 3, [[TMP24]]
+// CHECK: [[TMP26:%.*]] = and i64 [[TMP4]], 3
+// CHECK: [[TMP27:%.*]] = icmp eq i64 [[TMP26]], 0
+// CHECK: br i1 [[TMP27]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]]
// CHECK: omp.type.alloc:
-// CHECK: [[TMP36:%.*]] = and i64 [[TMP33]], -4
+// CHECK: [[TMP28:%.*]] = and i64 [[TMP25]], -4
// CHECK: br label [[OMP_TYPE_END:%.*]]
// CHECK: omp.type.alloc.else:
-// CHECK: [[TMP37:%.*]] = icmp eq i64 [[TMP34]], 1
-// CHECK: br i1 [[TMP37]], label [[OMP_TYPE_TO:%.*]], label [[OMP_TYPE_TO_ELSE:%.*]]
+// CHECK: [[TMP29:%.*]] = icmp eq i64 [[TMP26]], 1
+// CHECK: br i1 [[TMP29]], label [[OMP_TYPE_TO:%.*]], label [[OMP_TYPE_TO_ELSE:%.*]]
// CHECK: omp.type.to:
-// CHECK: [[TMP38:%.*]] = and i64 [[TMP33]], -3
+// CHECK: [[TMP30:%.*]] = and i64 [[TMP25]], -3
// CHECK: br label [[OMP_TYPE_END]]
// CHECK: omp.type.to.else:
-// CHECK: [[TMP39:%.*]] = icmp eq i64 [[TMP34]], 2
-// CHECK: br i1 [[TMP39]], label [[OMP_TYPE_FROM:%.*]], label [[OMP_TYPE_END]]
+// CHECK: [[TMP31:%.*]] = icmp eq i64 [[TMP26]], 2
+// CHECK: br i1 [[TMP31]], label [[OMP_TYPE_FROM:%.*]], label [[OMP_TYPE_END]]
// CHECK: omp.type.from:
-// CHECK: [[TMP40:%.*]] = and i64 [[TMP33]], -2
+// CHECK: [[TMP32:%.*]] = and i64 [[TMP25]], -2
// CHECK: br label [[OMP_TYPE_END]]
// CHECK: omp.type.end:
-// CHECK: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP36]], [[OMP_TYPE_ALLOC]] ], [ [[TMP38]], [[OMP_TYPE_TO]] ], [ [[TMP40]], [[OMP_TYPE_FROM]] ], [ [[TMP33]], [[OMP_TYPE_TO_ELSE]] ]
-// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP7]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[S1P]], i64 [[TMP30]], i64 [[OMP_MAPTYPE]], ptr null)
-// CHECK: [[TMP41:%.*]] = add nuw i64 281474976710659, [[TMP32]]
-// CHECK: [[TMP42:%.*]] = and i64 [[TMP12]], 3
-// CHECK: [[TMP43:%.*]] = icmp eq i64 [[TMP42]], 0
-// CHECK: br i1 [[TMP43]], label [[OMP_TYPE_ALLOC9:%.*]], label [[OMP_TYPE_ALLOC_ELSE10:%.*]]
-// CHECK: omp.type.alloc9:
-// CHECK: [[TMP44:%.*]] = and i64 [[TMP41]], -4
-// CHECK: br label [[OMP_TYPE_END14:%.*]]
-// CHECK: omp.type.alloc.else10:
-// CHECK: [[TMP45:%.*]] = icmp eq i64 [[TMP42]], 1
-// CHECK: br i1 [[TMP45]], label [[OMP_TYPE_TO11:%.*]], label [[OMP_TYPE_TO_ELSE12:%.*]]
-// CHECK: omp.type.to11:
-// CHECK: [[TMP46:%.*]] = and i64 [[TMP41]], -3
-// CHECK: br label [[OMP_TYPE_END14]]
-// CHECK: omp.type.to.else12:
-// CHECK: [[TMP47:%.*]] = icmp eq i64 [[TMP42]], 2
-// CHECK: br i1 [[TMP47]], label [[OMP_TYPE_FROM13:%.*]], label [[OMP_TYPE_END14]]
-// CHECK: omp.type.from13:
-// CHECK: [[TMP48:%.*]] = and i64 [[TMP41]], -2
-// CHECK: br label [[OMP_TYPE_END14]]
-// CHECK: omp.type.end14:
-// CHECK: [[OMP_MAPTYPE15:%.*]] = phi i64 [ [[TMP44]], [[OMP_TYPE_ALLOC9]] ], [ [[TMP46]], [[OMP_TYPE_TO11]] ], [ [[TMP48]], [[OMP_TYPE_FROM13]] ], [ [[TMP41]], [[OMP_TYPE_TO_ELSE12]] ]
-// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP7]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[Z]], i64 4, i64 [[OMP_MAPTYPE15]], ptr null)
-// CHECK: [[TMP49:%.*]] = add nuw i64 281474976710675, [[TMP32]]
-// CHECK: [[TMP50:%.*]] = and i64 [[TMP12]], 3
-// CHECK: [[TMP51:%.*]] = icmp eq i64 [[TMP50]], 0
-// CHECK: br i1 [[TMP51]], label [[OMP_TYPE_ALLOC16:%.*]], label [[OMP_TYPE_ALLOC_ELSE17:%.*]]
-// CHECK: omp.type.alloc16:
-// CHECK: [[TMP52:%.*]] = and i64 [[TMP49]], -4
-// CHECK: br label [[OMP_TYPE_END21:%.*]]
-// CHECK: omp.type.alloc.else17:
-// CHECK: [[TMP53:%.*]] = icmp eq i64 [[TMP50]], 1
-// CHECK: br i1 [[TMP53]], label [[OMP_TYPE_TO18:%.*]], label [[OMP_TYPE_TO_ELSE19:%.*]]
-// CHECK: omp.type.to18:
-// CHECK: [[TMP54:%.*]] = and i64 [[TMP49]], -3
-// CHECK: br label [[OMP_TYPE_END21]]
-// CHECK: omp.type.to.else19:
-// CHECK: [[TMP55:%.*]] = icmp eq i64 [[TMP50]], 2
-// CHECK: br i1 [[TMP55]], label [[OMP_TYPE_FROM20:%.*]], label [[OMP_TYPE_END21]]
-// CHECK: omp.type.from20:
-// CHECK: [[TMP56:%.*]] = and i64 [[TMP49]], -2
-// CHECK: br label [[OMP_TYPE_END21]]
-// CHECK: omp.type.end21:
-// CHECK: [[OMP_MAPTYPE22:%.*]] = phi i64 [ [[TMP52]], [[OMP_TYPE_ALLOC16]] ], [ [[TMP54]], [[OMP_TYPE_TO18]] ], [ [[TMP56]], [[OMP_TYPE_FROM20]] ], [ [[TMP49]], [[OMP_TYPE_TO_ELSE19]] ]
-// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP7]], ptr [[S1P]], ptr [[X]], i64 4, i64 [[OMP_MAPTYPE22]], ptr null)
-// CHECK: [[TMP57:%.*]] = add nuw i64 281474976710675, [[TMP32]]
-// CHECK: [[TMP58:%.*]] = and i64 [[TMP12]], 3
-// CHECK: [[TMP59:%.*]] = icmp eq i64 [[TMP58]], 0
-// CHECK: br i1 [[TMP59]], label [[OMP_TYPE_ALLOC23:%.*]], label [[OMP_TYPE_ALLOC_ELSE24:%.*]]
-// CHECK: omp.type.alloc23:
-// CHECK: [[TMP60:%.*]] = and i64 [[TMP57]], -4
-// CHECK: br label [[OMP_TYPE_END28]]
-// CHECK: omp.type.alloc.else24:
-// CHECK: [[TMP61:%.*]] = icmp eq i64 [[TMP58]], 1
-// CHECK: br i1 [[TMP61]], label [[OMP_TYPE_TO25:%.*]], label [[OMP_TYPE_TO_ELSE26:%.*]]
-// CHECK: omp.type.to25:
-// CHECK: [[TMP62:%.*]] = and i64 [[TMP57]], -3
-// CHECK: br label [[OMP_TYPE_END28]]
-// CHECK: omp.type.to.else26:
-// CHECK: [[TMP63:%.*]] = icmp eq i64 [[TMP58]], 2
-// CHECK: br i1 [[TMP63]], label [[OMP_TYPE_FROM27:%.*]], label [[OMP_TYPE_END28]]
-// CHECK: omp.type.from27:
-// CHECK: [[TMP64:%.*]] = and i64 [[TMP57]], -2
-// CHECK: br label [[OMP_TYPE_END28]]
-// CHECK: omp.type.end28:
-// CHECK: [[OMP_MAPTYPE29:%.*]] = phi i64 [ [[TMP60]], [[OMP_TYPE_ALLOC23]] ], [ [[TMP62]], [[OMP_TYPE_TO25]] ], [ [[TMP64]], [[OMP_TYPE_FROM27]] ], [ [[TMP57]], [[OMP_TYPE_TO_ELSE26]] ]
-// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP7]], ptr [[S1P7]], ptr [[Y]], i64 4, i64 [[OMP_MAPTYPE29]], ptr null)
+// CHECK: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP28]], [[OMP_TYPE_ALLOC]] ], [ [[TMP30]], [[OMP_TYPE_TO]] ], [ [[TMP32]], [[OMP_TYPE_FROM]] ], [ [[TMP25]], [[OMP_TYPE_TO_ELSE]] ]
+// CHECK: [[TMP33:%.*]] = and i64 [[TMP4]], 5132
+// CHECK: [[OMP_MAPTYPE_WITH_MODIFIERS:%.*]] = or i64 [[OMP_MAPTYPE]], [[TMP33]]
+// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[Z]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS]], ptr null)
+// CHECK: [[TMP34:%.*]] = and i64 [[TMP4]], 3
+// CHECK: [[TMP35:%.*]] = icmp eq i64 [[TMP34]], 0
+// CHECK: br i1 [[TMP35]], label [[OMP_TYPE_ALLOC4:%.*]], label [[OMP_TYPE_ALLOC_ELSE5:%.*]]
+// CHECK: omp.type.alloc4:
+// CHECK: br label [[OMP_TYPE_END9:%.*]]
+// CHECK: omp.type.alloc.else5:
+// CHECK: [[TMP36:%.*]] = icmp eq i64 [[TMP34]], 1
+// CHECK: br i1 [[TMP36]], label [[OMP_TYPE_TO6:%.*]], label [[OMP_TYPE_TO_ELSE7:%.*]]
+// CHECK: omp.type.to6:
+// CHECK: br label [[OMP_TYPE_END9]]
+// CHECK: omp.type.to.else7:
+// CHECK: [[TMP37:%.*]] = icmp eq i64 [[TMP34]], 2
+// CHECK: br i1 [[TMP37]], label [[OMP_TYPE_FROM8:%.*]], label [[OMP_TYPE_END9]]
+// CHECK: omp.type.from8:
+// CHECK: br label [[OMP_TYPE_END9]]
+// CHECK: omp.type.end9:
+// CHECK: [[OMP_MAPTYPE10:%.*]] = phi i64 [ 0, [[OMP_TYPE_ALLOC4]] ], [ 0, [[OMP_TYPE_TO6]] ], [ 0, [[OMP_TYPE_FROM8]] ], [ 0, [[OMP_TYPE_TO_ELSE7]] ]
+// CHECK: [[TMP38:%.*]] = and i64 [[TMP4]], 5132
+// CHECK: [[OMP_MAPTYPE_WITH_MODIFIERS11:%.*]] = or i64 [[OMP_MAPTYPE10]], [[TMP38]]
+// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP15]], ptr [[X]], i64 [[TMP22]], i64 [[OMP_MAPTYPE_WITH_MODIFIERS11]], ptr null)
+// CHECK: [[TMP39:%.*]] = add nuw i64 562949953421315, [[TMP24]]
+// CHECK: [[TMP40:%.*]] = and i64 [[TMP4]], 3
+// CHECK: [[TMP41:%.*]] = icmp eq i64 [[TMP40]], 0
+// CHECK: br i1 [[TMP41]], label [[OMP_TYPE_ALLOC12:%.*]], label [[OMP_TYPE_ALLOC_ELSE13:%.*]]
+// CHECK: omp.type.alloc12:
+// CHECK: [[TMP42:%.*]] = and i64 [[TMP39]], -4
+// CHECK: br label [[OMP_TYPE_END17:%.*]]
+// CHECK: omp.type.alloc.else13:
+// CHECK: [[TMP43:%.*]] = icmp eq i64 [[TMP40]], 1
+// CHECK: br i1 [[TMP43]], label [[OMP_TYPE_TO14:%.*]], label [[OMP_TYPE_TO_ELSE15:%.*]]
+// CHECK: omp.type.to14:
+// CHECK: [[TMP44:%.*]] = and i64 [[TMP39]], -3
+// CHECK: br label [[OMP_TYPE_END17]]
+// CHECK: omp.type.to.else15:
+// CHECK: [[TMP45:%.*]] = icmp eq i64 [[TMP40]], 2
+// CHECK: br i1 [[TMP45]], label [[OMP_TYPE_FROM16:%.*]], label [[OMP_TYPE_END17]]
+// CHECK: omp.type.from16:
+// CHECK: [[TMP46:%.*]] = and i64 [[TMP39]], -2
+// CHECK: br label [[OMP_TYPE_END17]]
+// CHECK: omp.type.end17:
+// CHECK: [[OMP_MAPTYPE18:%.*]] = phi i64 [ [[TMP42]], [[OMP_TYPE_ALLOC12]] ], [ [[TMP44]], [[OMP_TYPE_TO14]] ], [ [[TMP46]], [[OMP_TYPE_FROM16]] ], [ [[TMP39]], [[OMP_TYPE_TO_ELSE15]] ]
+// CHECK: [[TMP47:%.*]] = and i64 [[TMP4]], 5132
+// CHECK: [[OMP_MAPTYPE_WITH_MODIFIERS19:%.*]] = or i64 [[OMP_MAPTYPE18]], [[TMP47]]
+// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP15]], ptr [[X]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS19]], ptr null)
+// CHECK: [[TMP48:%.*]] = add nuw i64 562949953421315, [[TMP24]]
+// CHECK: [[TMP49:%.*]] = and i64 [[TMP4]], 3
+// CHECK: [[TMP50:%.*]] = icmp eq i64 [[TMP49]], 0
+// CHECK: br i1 [[TMP50]], label [[OMP_TYPE_ALLOC20:%.*]], label [[OMP_TYPE_ALLOC_ELSE21:%.*]]
+// CHECK: omp.type.alloc20:
+// CHECK: [[TMP51:%.*]] = and i64 [[TMP48]], -4
+// CHECK: br label [[OMP_TYPE_END25:%.*]]
+// CHECK: omp.type.alloc.else21:
+// CHECK: [[TMP52:%.*]] = icmp eq i64 [[TMP49]], 1
+// CHECK: br i1 [[TMP52]], label [[OMP_TYPE_TO22:%.*]], label [[OMP_TYPE_TO_ELSE23:%.*]]
+// CHECK: omp.type.to22:
+// CHECK: [[TMP53:%.*]] = and i64 [[TMP48]], -3
+// CHECK: br label [[OMP_TYPE_END25]]
+// CHECK: omp.type.to.else23:
+// CHECK: [[TMP54:%.*]] = icmp eq i64 [[TMP49]], 2
+// CHECK: br i1 [[TMP54]], label [[OMP_TYPE_FROM24:%.*]], label [[OMP_TYPE_END25]]
+// CHECK: omp.type.from24:
+// CHECK: [[TMP55:%.*]] = and i64 [[TMP48]], -2
+// CHECK: br label [[OMP_TYPE_END25]]
+// CHECK: omp.type.end25:
+// CHECK: [[OMP_MAPTYPE26:%.*]] = phi i64 [ [[TMP51]], [[OMP_TYPE_ALLOC20]] ], [ [[TMP53]], [[OMP_TYPE_TO22]] ], [ [[TMP55]], [[OMP_TYPE_FROM24]] ], [ [[TMP48]], [[OMP_TYPE_TO_ELSE23]] ]
+// CHECK: [[TMP56:%.*]] = and i64 [[TMP4]], 5132
+// CHECK: [[OMP_MAPTYPE_WITH_MODIFIERS27:%.*]] = or i64 [[OMP_MAPTYPE26]], [[TMP56]]
+// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP17]], ptr [[Y]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS27]], ptr null)
+// CHECK: [[TMP57:%.*]] = and i64 [[TMP4]], 3
+// CHECK: [[TMP58:%.*]] = icmp eq i64 [[TMP57]], 0
+// CHECK: br i1 [[TMP58]], label [[OMP_TYPE_ALLOC28:%.*]], label [[OMP_TYPE_ALLOC_ELSE29:%.*]]
+// CHECK: omp.type.alloc28:
+// CHECK: br label [[OMP_TYPE_END33]]
+// CHECK: omp.type.alloc.else29:
+// CHECK: [[TMP59:%.*]] = icmp eq i64 [[TMP57]], 1
+// CHECK: br i1 [[TMP59]], label [[OMP_TYPE_TO30:%.*]], label [[OMP_TYPE_TO_ELSE31:%.*]]
+// CHECK: omp.type.to30:
+// CHECK: br label [[OMP_TYPE_END33]]
+// CHECK: omp.type.to.else31:
+// CHECK: [[TMP60:%.*]] = icmp eq i64 [[TMP57]], 2
+// CHECK: br i1 [[TMP60]], label [[OMP_TYPE_FROM32:%.*]], label [[OMP_TYPE_END33]]
+// CHECK: omp.type.from32:
+// CHECK: br label [[OMP_TYPE_END33]]
+// CHECK: omp.type.end33:
+// CHECK: [[OMP_MAPTYPE34:%.*]] = phi i64 [ 16384, [[OMP_TYPE_ALLOC28]] ], [ 16384, [[OMP_TYPE_TO30]] ], [ 16384, [[OMP_TYPE_FROM32]] ], [ 16384, [[OMP_TYPE_TO_ELSE31]] ]
+// CHECK: [[TMP61:%.*]] = and i64 [[TMP4]], 5132
+// CHECK: [[OMP_MAPTYPE_WITH_MODIFIERS35:%.*]] = or i64 [[OMP_MAPTYPE34]], [[TMP61]]
+// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[S1P2]], ptr [[X]], i64 8, i64 [[OMP_MAPTYPE34]], ptr null)
// CHECK: [[OMP_ARRAYMAP_NEXT]] = getelementptr [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 1
-// CHECK: [[OMP_ARRAYMAP_ISDONE:%.*]] = icmp eq ptr [[OMP_ARRAYMAP_NEXT]], [[TMP11]]
+// CHECK: [[OMP_ARRAYMAP_ISDONE:%.*]] = icmp eq ptr [[OMP_ARRAYMAP_NEXT]], [[TMP7]]
// CHECK: br i1 [[OMP_ARRAYMAP_ISDONE]], label [[OMP_ARRAYMAP_EXIT:%.*]], label [[OMP_ARRAYMAP_BODY]]
// CHECK: omp.arraymap.exit:
-// CHECK: [[OMP_ARRAYINIT_ISARRAY30:%.*]] = icmp sgt i64 [[TMP10]], 1
-// CHECK: [[TMP65:%.*]] = and i64 [[TMP12]], 8
-// CHECK: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP65]], 0
-// CHECK: [[TMP66:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY30]], [[DOTOMP_ARRAY__DEL__DELETE]]
-// CHECK: br i1 [[TMP66]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]]
+// CHECK: [[OMP_ARRAYINIT_ISARRAY36:%.*]] = icmp sgt i64 [[TMP6]], 1
+// CHECK: [[TMP62:%.*]] = and i64 [[TMP4]], 8
+// CHECK: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP62]], 0
+// CHECK: [[TMP63:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY36]], [[DOTOMP_ARRAY__DEL__DELETE]]
+// CHECK: br i1 [[TMP63]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]]
// CHECK: .omp.array..del:
-// CHECK: [[TMP67:%.*]] = mul nuw i64 [[TMP10]], 16
-// CHECK: [[TMP68:%.*]] = and i64 [[TMP12]], -4
-// CHECK: [[TMP69:%.*]] = or i64 [[TMP68]], 512
-// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP7]], ptr [[TMP8]], ptr [[TMP9]], i64 [[TMP67]], i64 [[TMP69]], ptr [[TMP13]])
+// CHECK: [[TMP64:%.*]] = mul nuw i64 [[TMP6]], 16
+// CHECK: [[TMP65:%.*]] = and i64 [[TMP4]], -4
+// CHECK: [[TMP66:%.*]] = or i64 [[TMP65]], 512
+// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP64]], i64 [[TMP66]], ptr [[TMP5]])
// CHECK: br label [[OMP_DONE]]
// CHECK: omp.done:
// CHECK: ret void
diff --git a/offload/test/mapping/declare_mapper_nested_mappers.cpp b/offload/test/mapping/declare_mapper_nested_mappers.cpp
index a59ed6980ec4c..5966bb6127af8 100644
--- a/offload/test/mapping/declare_mapper_nested_mappers.cpp
+++ b/offload/test/mapping/declare_mapper_nested_mappers.cpp
@@ -7,7 +7,8 @@ typedef struct {
int a;
double *b;
} C;
-#pragma omp declare mapper(id1 : C s) map(to : s.a) map(from : s.b[0 : 2])
+#pragma omp declare mapper(id1 : C s) map(to : s.a) map(alloc : s.b) \
+ map(from : s.b[0 : 2])
typedef struct {
int e;
@@ -16,7 +17,7 @@ typedef struct {
short *g;
} D;
#pragma omp declare mapper(default : D r) map(from : r.e) \
- map(mapper(id1), tofrom : r.f) map(tofrom : r.g[0 : r.h])
+ map(mapper(id1), tofrom : r.f) map(alloc : r.g) map(tofrom : r.g[0 : r.h])
int main() {
constexpr int N = 10;
@@ -56,7 +57,7 @@ int main() {
spp[0][0].f.b[1] = 40;
spp[0][0].g[1] = 50;
}
- printf("%d %d %d %d\n", spp00fa, spp00fb_r, spp00fg1, spp00fg_r);
+ printf("%d %d %d %d\n", spp00fa, spp00fb_r, spp00fg1, spp00fg_r);
// CHECK: 222 0 30 0
printf("%d %d %4.5f %d %d %d\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.b[1],
spp[0][0].f.b == &x[0] ? 1 : 0, spp[0][0].g[1],
diff --git a/offload/test/mapping/declare_mapper_target.cpp b/offload/test/mapping/declare_mapper_target.cpp
index 4d7237e94657a..c642949b2a367 100644
--- a/offload/test/mapping/declare_mapper_target.cpp
+++ b/offload/test/mapping/declare_mapper_target.cpp
@@ -10,7 +10,7 @@ class C {
int *a;
};
-#pragma omp declare mapper(id : C s) map(s.a[0 : NUM])
+#pragma omp declare mapper(id : C s) map(alloc : s.a) map(s.a[0 : NUM])
int main() {
C c;
diff --git a/offload/test/mapping/declare_mapper_target_data.cpp b/offload/test/mapping/declare_mapper_target_data.cpp
index 7f09844014005..fd2c8f4886401 100644
--- a/offload/test/mapping/declare_mapper_target_data.cpp
+++ b/offload/test/mapping/declare_mapper_target_data.cpp
@@ -10,7 +10,7 @@ class C {
int *a;
};
-#pragma omp declare mapper(id : C s) map(s.a[0 : NUM])
+#pragma omp declare mapper(id : C s) map(alloc : s.a) map(s.a[0 : NUM])
int main() {
C c;
diff --git a/offload/test/mapping/declare_mapper_target_data_enter_exit.cpp b/offload/test/mapping/declare_mapper_target_data_enter_exit.cpp
index f5fad8b8fe332..245462eab5785 100644
--- a/offload/test/mapping/declare_mapper_target_data_enter_exit.cpp
+++ b/offload/test/mapping/declare_mapper_target_data_enter_exit.cpp
@@ -10,7 +10,7 @@ class C {
int *a;
};
-#pragma omp declare mapper(id : C s) map(s.a[0 : NUM])
+#pragma omp declare mapper(id : C s) map(alloc : s.a) map(s.a[0 : NUM])
int main() {
C c;
diff --git a/offload/test/mapping/declare_mapper_target_update.cpp b/offload/test/mapping/declare_mapper_target_update.cpp
index fe4597b76908f..4e053e3e56139 100644
--- a/offload/test/mapping/declare_mapper_target_update.cpp
+++ b/offload/test/mapping/declare_mapper_target_update.cpp
@@ -10,7 +10,7 @@ class C {
int *a;
};
-#pragma omp declare mapper(id : C s) map(s.a[0 : NUM])
+#pragma omp declare mapper(id : C s) map(alloc : s.a) map(s.a[0 : NUM])
int main() {
C c;
>From a59048f249d20037fa3a0151cb1d0ce6ee9a96a7 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Fri, 24 Apr 2026 03:41:59 -0700
Subject: [PATCH 09/13] Add one more test.
---
.../mapper_map_ptee_only_always_array.c | 70 +++++++++++++++++++
1 file changed, 70 insertions(+)
create mode 100644 offload/test/mapping/mapper_map_ptee_only_always_array.c
diff --git a/offload/test/mapping/mapper_map_ptee_only_always_array.c b/offload/test/mapping/mapper_map_ptee_only_always_array.c
new file mode 100644
index 0000000000000..48e2a492131d3
--- /dev/null
+++ b/offload/test/mapping/mapper_map_ptee_only_always_array.c
@@ -0,0 +1,70 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+// Test that a mapper that maps a var and a pointee, correctly propagates the
+// always and present bits into the individual maps "pushed" by it.
+
+int x[2][10];
+
+typedef struct {
+ int x;
+ int dummy[10000];
+ int *p;
+} S;
+
+#pragma omp declare mapper(default : S s) map(s.x, s.p[0 : 10])
+
+S s1[2];
+
+void print_status(void *p, const char *name) {
+ int present = omp_target_is_present(p, omp_get_default_device());
+ printf("%s is %spresent\n", name, present ? "" : "not ");
+}
+
+int main() {
+ x[0][0] = x[0][1] = s1[0].x = s1[1].x = 111;
+ s1[0].p = (int *)&x[0];
+ s1[1].p = (int *)&x[1];
+
+#pragma omp target enter data map(alloc : s1)
+ printf("After mapping\n");
+ print_status(&s1[0].x, "s1[0].x"); // CHECK: s1[0].x is present
+ // dummy/p being present is not ideal, but that's what we get with the
+ // current implementation because we need to map the full contiguous
+ // chunk for the array first before invoking the mapper.
+ print_status(&s1[0].dummy, "s1[0].dummy"); // CHECK: s1[0].dummy is present
+ print_status(&s1[0].p, "s1[0].p"); // CHECK: s1[0].p is present
+ print_status(&s1[0].p[0], "s1[0].p[0]"); // CHECK: s1[0].p[0] is present
+ print_status(&s1[1].x, "s1[1].x"); // CHECK: s1[1].x is present
+ print_status(&s1[1].dummy, "s1[1].dummy"); // CHECK: s1[1].dummy is present
+ print_status(&s1[1].p, "s1[1].p"); // CHECK: s1[1].p is present
+ print_status(&s1[1].p[0], "s1[1].p[0]"); // CHECK: s1[1].p[0] is present
+
+#pragma omp target map(always, present, from : s1)
+ {
+ s1[0].p[0] = s1[1].p[0] = s1[0].x = s1[1].x = 222;
+ }
+
+ printf("\n");
+ printf("After map(always,from)\n");
+ printf("s[0].x = %d\n", s1[0].x); // CHECK: s[0].x = 222
+ printf("s[1].x = %d\n", s1[1].x); // CHECK: s[1].x = 222
+ printf("s[0].p[0] = %d\n", s1[0].p[0]); // CHECK: s[0].p[0] = 222
+ printf("s[1].p[0] = %d\n", s1[1].p[0]); // CHECK: s[1].p[0] = 222
+ printf("\n");
+
+#pragma omp target exit data map(delete : s1)
+ printf("After deleting\n");
+ print_status(&s1[0].x, "s1[0].x"); // CHECK: s1[0].x is not present
+ print_status(&s1[0].dummy,
+ "s1[0].dummy"); // CHECK: s1[0].dummy is not present
+ print_status(&s1[0].p, "s1[0].p"); // CHECK: s1[0].p is not present
+ print_status(&s1[0].p[0], "s1[0].p[0]"); // CHECK: s1[0].p[0] is not present
+ print_status(&s1[1].x, "s1[1].x"); // CHECK: s1[1].x is not present
+ print_status(&s1[1].dummy,
+ "s1[1].dummy"); // CHECK: s1[1].dummy is not present
+ print_status(&s1[1].p, "s1[1].p"); // CHECK: s1[1].p is not present
+ print_status(&s1[1].p[0], "s1[1].p[0]"); // CHECK: s1[1].p[0] is not present
+}
>From 046c264327180ea23e2f13a4c4a8618704815fca Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Fri, 24 Apr 2026 03:52:31 -0700
Subject: [PATCH 10/13] Clang-format fixes.
---
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 6 +-
...apper_map_mbr_ptee_then_present_mbr_ptee.c | 4 +-
...r_map_ptee_only_2_ptr_indirections_array.c | 66 ++++++++++++-------
.../mapper_map_ptee_only_2ndlevel_array.c | 6 +-
4 files changed, 50 insertions(+), 32 deletions(-)
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 993456b3db28c..e7801517043f9 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -10118,15 +10118,15 @@ Expected<Function *> OpenMPIRBuilder::emitUserDefinedMapper(
CurMapType, ModifierBitMask, "omp.maptype.with.modifiers");
// ATTACH entries must not receive map-type-modifying bits: ATTACH|ALWAYS is
- // reserved for the attach(always) map-type modifier, and other modifier bits
- // (DELETE, CLOSE, PRESENT) have no meaning for an ATTACH entry.
+ // reserved for the attach(always) map-type modifier, and other modifier
+ // bits (DELETE, CLOSE, PRESENT) have no meaning for an ATTACH entry.
constexpr uint64_t AttachBit =
static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>(
OpenMPOffloadMappingFlags::OMP_MAP_ATTACH);
Value *FinalMapType =
(RawType & AttachBit) ? CurMapType : CurMapTypeWithModifiers;
- Value *OffloadingArgs[] = {MapperHandle, CurBaseArg, CurBeginArg,
+ Value *OffloadingArgs[] = {MapperHandle, CurBaseArg, CurBeginArg,
CurSizeArg, FinalMapType, CurNameArg};
auto ChildMapperFn = CustomMapperCB(I);
diff --git a/offload/test/mapping/mapper_map_mbr_ptee_then_present_mbr_ptee.c b/offload/test/mapping/mapper_map_mbr_ptee_then_present_mbr_ptee.c
index 00a8ba4bb4a81..8fbba33136521 100644
--- a/offload/test/mapping/mapper_map_mbr_ptee_then_present_mbr_ptee.c
+++ b/offload/test/mapping/mapper_map_mbr_ptee_then_present_mbr_ptee.c
@@ -26,7 +26,7 @@ void print_status(void *p, const char *name) {
int main() {
s1.p = (int *)&x;
-#pragma omp target enter data map(alloc : s1.x, s1.p[0:10])
+#pragma omp target enter data map(alloc : s1.x, s1.p[0 : 10])
printf("After mapping\n");
print_status(&s1.x, "x"); // CHECK: x is present
print_status(&s1.dummy, "dummy"); // CHECK: dummy is not present
@@ -37,7 +37,7 @@ int main() {
// This present check should pass.
#pragma omp target enter data map(present, alloc : s1)
-#pragma omp target exit data map(delete: s1)
+#pragma omp target exit data map(delete : s1)
printf("After deleting\n");
print_status(&s1.x, "x"); // CHECK: x is not present
print_status(&s1.dummy, "dummy"); // CHECK: dummy is not present
diff --git a/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections_array.c b/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections_array.c
index 621b78ba17061..076867f05470e 100644
--- a/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections_array.c
+++ b/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections_array.c
@@ -23,7 +23,7 @@ typedef struct {
} S2;
#pragma omp declare mapper(default : S2 s2) \
- map(s2.z, s2.s1p->x, s2.s1p->y, s2.s1p->p[0 : 10])
+ map(s2.z, s2.s1p -> x, s2.s1p->y, s2.s1p->p[0 : 10])
S1 s1arr[2];
S2 s2arr[2];
@@ -41,32 +41,48 @@ int main() {
#pragma omp target enter data map(alloc : s2arr)
printf("After mapping\n");
- print_status(&s2arr[0].s1p->x, "s2arr[0].x"); // CHECK: s2arr[0].x is present
- print_status(&s2arr[0].s1p->y, "s2arr[0].y"); // CHECK: s2arr[0].y is present
- print_status(&s2arr[0].z, "s2arr[0].z"); // CHECK: s2arr[0].z is present
- print_status(&s2arr[0].s1p->dummy, "s2arr[0].dummy"); // CHECK: s2arr[0].dummy is not present
- print_status(&s2arr[0].s1p->p, "s2arr[0].p"); // CHECK: s2arr[0].p is not present
- print_status(&s2arr[0].s1p->p[0], "s2arr[0].p[0]"); // CHECK: s2arr[0].p[0] is present
- print_status(&s2arr[1].s1p->x, "s2arr[1].x"); // CHECK: s2arr[1].x is present
- print_status(&s2arr[1].s1p->y, "s2arr[1].y"); // CHECK: s2arr[1].y is present
- print_status(&s2arr[1].z, "s2arr[1].z"); // CHECK: s2arr[1].z is present
- print_status(&s2arr[1].s1p->dummy, "s2arr[1].dummy"); // CHECK: s2arr[1].dummy is not present
- print_status(&s2arr[1].s1p->p, "s2arr[1].p"); // CHECK: s2arr[1].p is not present
- print_status(&s2arr[1].s1p->p[0], "s2arr[1].p[0]"); // CHECK: s2arr[1].p[0] is present
+ print_status(&s2arr[0].s1p->x, "s2arr[0].x"); // CHECK: s2arr[0].x is present
+ print_status(&s2arr[0].s1p->y, "s2arr[0].y"); // CHECK: s2arr[0].y is present
+ print_status(&s2arr[0].z, "s2arr[0].z"); // CHECK: s2arr[0].z is present
+ print_status(&s2arr[0].s1p->dummy,
+ "s2arr[0].dummy"); // CHECK: s2arr[0].dummy is not present
+ print_status(&s2arr[0].s1p->p,
+ "s2arr[0].p"); // CHECK: s2arr[0].p is not present
+ print_status(&s2arr[0].s1p->p[0],
+ "s2arr[0].p[0]"); // CHECK: s2arr[0].p[0] is present
+ print_status(&s2arr[1].s1p->x, "s2arr[1].x"); // CHECK: s2arr[1].x is present
+ print_status(&s2arr[1].s1p->y, "s2arr[1].y"); // CHECK: s2arr[1].y is present
+ print_status(&s2arr[1].z, "s2arr[1].z"); // CHECK: s2arr[1].z is present
+ print_status(&s2arr[1].s1p->dummy,
+ "s2arr[1].dummy"); // CHECK: s2arr[1].dummy is not present
+ print_status(&s2arr[1].s1p->p,
+ "s2arr[1].p"); // CHECK: s2arr[1].p is not present
+ print_status(&s2arr[1].s1p->p[0],
+ "s2arr[1].p[0]"); // CHECK: s2arr[1].p[0] is present
printf("\n");
#pragma omp target exit data map(delete : s2arr)
printf("After deleting\n");
- print_status(&s2arr[0].s1p->x, "s2arr[0].x"); // CHECK: s2arr[0].x is not present
- print_status(&s2arr[0].s1p->y, "s2arr[0].y"); // CHECK: s2arr[0].y is not present
- print_status(&s2arr[0].z, "s2arr[0].z"); // CHECK: s2arr[0].z is not present
- print_status(&s2arr[0].s1p->dummy, "s2arr[0].dummy"); // CHECK: s2arr[0].dummy is not present
- print_status(&s2arr[0].s1p->p, "s2arr[0].p"); // CHECK: s2arr[0].p is not present
- print_status(&s2arr[0].s1p->p[0], "s2arr[0].p[0]"); // CHECK: s2arr[0].p[0] is not present
- print_status(&s2arr[1].s1p->x, "s2arr[1].x"); // CHECK: s2arr[1].x is not present
- print_status(&s2arr[1].s1p->y, "s2arr[1].y"); // CHECK: s2arr[1].y is not present
- print_status(&s2arr[1].z, "s2arr[1].z"); // CHECK: s2arr[1].z is not present
- print_status(&s2arr[1].s1p->dummy, "s2arr[1].dummy"); // CHECK: s2arr[1].dummy is not present
- print_status(&s2arr[1].s1p->p, "s2arr[1].p"); // CHECK: s2arr[1].p is not present
- print_status(&s2arr[1].s1p->p[0], "s2arr[1].p[0]"); // CHECK: s2arr[1].p[0] is not present
+ print_status(&s2arr[0].s1p->x,
+ "s2arr[0].x"); // CHECK: s2arr[0].x is not present
+ print_status(&s2arr[0].s1p->y,
+ "s2arr[0].y"); // CHECK: s2arr[0].y is not present
+ print_status(&s2arr[0].z, "s2arr[0].z"); // CHECK: s2arr[0].z is not present
+ print_status(&s2arr[0].s1p->dummy,
+ "s2arr[0].dummy"); // CHECK: s2arr[0].dummy is not present
+ print_status(&s2arr[0].s1p->p,
+ "s2arr[0].p"); // CHECK: s2arr[0].p is not present
+ print_status(&s2arr[0].s1p->p[0],
+ "s2arr[0].p[0]"); // CHECK: s2arr[0].p[0] is not present
+ print_status(&s2arr[1].s1p->x,
+ "s2arr[1].x"); // CHECK: s2arr[1].x is not present
+ print_status(&s2arr[1].s1p->y,
+ "s2arr[1].y"); // CHECK: s2arr[1].y is not present
+ print_status(&s2arr[1].z, "s2arr[1].z"); // CHECK: s2arr[1].z is not present
+ print_status(&s2arr[1].s1p->dummy,
+ "s2arr[1].dummy"); // CHECK: s2arr[1].dummy is not present
+ print_status(&s2arr[1].s1p->p,
+ "s2arr[1].p"); // CHECK: s2arr[1].p is not present
+ print_status(&s2arr[1].s1p->p[0],
+ "s2arr[1].p[0]"); // CHECK: s2arr[1].p[0] is not present
}
diff --git a/offload/test/mapping/mapper_map_ptee_only_2ndlevel_array.c b/offload/test/mapping/mapper_map_ptee_only_2ndlevel_array.c
index 6a7ee78c4544a..08b9551f0c8a9 100644
--- a/offload/test/mapping/mapper_map_ptee_only_2ndlevel_array.c
+++ b/offload/test/mapping/mapper_map_ptee_only_2ndlevel_array.c
@@ -36,7 +36,8 @@ int main() {
#pragma omp target enter data map(to : s2arr)
printf("After mapping\n");
- print_status(&s2arr[0].s1.x, "s2arr[0].s1.x"); // CHECK: s2arr[0].s1.x is present
+ print_status(&s2arr[0].s1.x,
+ "s2arr[0].s1.x"); // CHECK: s2arr[0].s1.x is present
// dummy/p being present is not ideal, but that's what we get with the
// current implementation because we need to map the full contiguous
// chunk for the array first before invoking the mapper.
@@ -46,7 +47,8 @@ int main() {
"s2arr[0].s1.p"); // CHECK: s2arr[0].s1.p is present
print_status(&s2arr[0].s1.p[0],
"s2arr[0].s1.p[0]"); // CHECK: s2arr[0].s1.p[0] is present
- print_status(&s2arr[1].s1.x, "s2arr[1].s1.x"); // CHECK: s2arr[1].s1.x is present
+ print_status(&s2arr[1].s1.x,
+ "s2arr[1].s1.x"); // CHECK: s2arr[1].s1.x is present
print_status(&s2arr[1].s1.dummy,
"s2arr[1].s1.dummy"); // CHECK: s2arr[1].s1.dummy is present
print_status(&s2arr[1].s1.p,
>From fbf40e9cf825bf794a5755a3e93ef564654fd0f1 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Wed, 13 May 2026 17:35:41 -0700
Subject: [PATCH 11/13] Minor comment/test changes.
---
..._of_structs_with_nested_mapper_codegen.cpp | 11 ++++++++++
..._of_structs_with_nested_mapper_codegen.cpp | 11 ++++++++++
...t_map_nested_ptr_member_mapper_codegen.cpp | 20 ++++++++++---------
mlir/test/Target/LLVMIR/omptarget-llvm.mlir | 4 +++-
.../mapper_map_mbr_then_present_mbr_ptee.c | 2 +-
.../mapper_map_ptee_only_2_ptr_indirections.c | 5 +++--
...r_map_ptee_only_2_ptr_indirections_array.c | 9 ++++++---
.../mapper_map_ptee_only_2ndlevel_array.c | 10 +++++++---
8 files changed, 53 insertions(+), 19 deletions(-)
diff --git a/clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_codegen.cpp b/clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_codegen.cpp
index 5703205167c8f..70514efe174af 100644
--- a/clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_codegen.cpp
+++ b/clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_codegen.cpp
@@ -117,6 +117,13 @@ void foo() {
// CHECK-NEXT: ret void
//
//
+// D mapper for: map(tofrom: sa) — per-element entries for struct D
+// (i = array element index, N = __tgt_mapper_num_components()):
+// &sa[i], &sa[i].e, sizeof(e..h), ALLOC | MEMBER_OF(N) | modifiers
+// &sa[i], &sa[i].e, sizeof(int), MEMBER_OF(N+1) | TO | FROM | modifiers
+// &sa[i], &sa[i].f, sizeof(C), MEMBER_OF(N+1) | TO | FROM | modifiers (dispatches to C mapper)
+// &sa[i], &sa[i].h, sizeof(int), MEMBER_OF(N+1) | TO | FROM | modifiers
+//
// CHECK-LABEL: define {{[^@]+}}@.omp_mapper._ZTS1D.default
// CHECK-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]], ptr noundef [[TMP2:%.*]], i64 noundef [[TMP3:%.*]], i64 noundef [[TMP4:%.*]], ptr noundef [[TMP5:%.*]]) #[[ATTR2:[0-9]+]] {
// CHECK-NEXT: entry:
@@ -264,6 +271,10 @@ void foo() {
// CHECK-NEXT: ret void
//
//
+// C mapper for: map(to: s.a) — per-element entries for struct C
+// (i = array element index, N = __tgt_mapper_num_components()):
+// &c[i], &c[i].a, sizeof(int), MEMBER_OF(N) | TO | FROM | modifiers
+//
// CHECK-LABEL: define {{[^@]+}}@.omp_mapper._ZTS1C.default
// CHECK-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]], ptr noundef [[TMP2:%.*]], i64 noundef [[TMP3:%.*]], i64 noundef [[TMP4:%.*]], ptr noundef [[TMP5:%.*]]) #[[ATTR2]] {
// CHECK-NEXT: entry:
diff --git a/clang/test/OpenMP/target_map_array_section_of_structs_with_nested_mapper_codegen.cpp b/clang/test/OpenMP/target_map_array_section_of_structs_with_nested_mapper_codegen.cpp
index f72a2e68c0d3e..fdc2eb31169ef 100644
--- a/clang/test/OpenMP/target_map_array_section_of_structs_with_nested_mapper_codegen.cpp
+++ b/clang/test/OpenMP/target_map_array_section_of_structs_with_nested_mapper_codegen.cpp
@@ -114,6 +114,13 @@ void foo() {
// CHECK-NEXT: ret void
//
//
+// D mapper for: map(tofrom: sa[0:2]) — per-element entries for struct D
+// (i = array element index, N = __tgt_mapper_num_components()):
+// &sa[i], &sa[i].e, sizeof(e..h), ALLOC | MEMBER_OF(N) | modifiers
+// &sa[i], &sa[i].e, sizeof(int), MEMBER_OF(N+1) | TO | FROM | modifiers
+// &sa[i], &sa[i].f, sizeof(C), MEMBER_OF(N+1) | TO | FROM | modifiers (dispatches to C mapper)
+// &sa[i], &sa[i].h, sizeof(int), MEMBER_OF(N+1) | TO | FROM | modifiers
+//
// CHECK-LABEL: define {{[^@]+}}@.omp_mapper._ZTS1D.default
// CHECK-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]], ptr noundef [[TMP2:%.*]], i64 noundef [[TMP3:%.*]], i64 noundef [[TMP4:%.*]], ptr noundef [[TMP5:%.*]]) #[[ATTR2:[0-9]+]] {
// CHECK-NEXT: entry:
@@ -261,6 +268,10 @@ void foo() {
// CHECK-NEXT: ret void
//
//
+// C mapper for: map(to: s.a) — per-element entries for struct C
+// (i = array element index, N = __tgt_mapper_num_components()):
+// &c[i], &c[i].a, sizeof(int), MEMBER_OF(N) | TO | modifiers
+//
// CHECK-LABEL: define {{[^@]+}}@.omp_mapper._ZTS1C.default
// CHECK-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]], ptr noundef [[TMP2:%.*]], i64 noundef [[TMP3:%.*]], i64 noundef [[TMP4:%.*]], ptr noundef [[TMP5:%.*]]) #[[ATTR2]] {
// CHECK-NEXT: entry:
diff --git a/clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp b/clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp
index 82509ace56eec..320ccabeea5c8 100644
--- a/clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp
+++ b/clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp
@@ -7,16 +7,18 @@
#ifndef HEADER
#define HEADER
-// &arr[i], &arr[i].z, sizeof(int), MEMBER_OF(N) | TO | FROM
-// &arr[i].s1p[0], &arr[i].s1p->x, sizeof(s1p->x..y), ALLOC (*)
-// &arr[i].s1p[0], &arr[i].s1p->x, sizeof(int), MEMBER_OF(N+2) | TO | FROM (**)
-// &arr[i].s1p[0], &arr[i].s1p->y, sizeof(int), MEMBER_OF(N+2) | TO | FROM (**)
+// S2 mapper for: map(to: arr[0:2])
+// Per-element entries (i = array element index; N = __tgt_mapper_num_components()):
+// &arr[i], &arr[i].z, sizeof(int), MEMBER_OF(N) | TO | FROM | modifiers
+// &arr[i].s1p[0], &arr[i].s1p->x, sizeof(s1p->x..y), ALLOC | modifiers (*)
+// &arr[i].s1p[0], &arr[i].s1p->x, sizeof(int), MEMBER_OF(N+2) | TO | FROM | modifiers (**)
+// &arr[i].s1p[0], &arr[i].s1p->y, sizeof(int), MEMBER_OF(N+2) | TO | FROM | modifiers (**)
// &arr[i].s1p, &arr[i].s1p->x, sizeof(ptr), ATTACH (***)
-// (*) Pointee (combined entry): No MEMBER_OF addition
-// (**) Inner MEMBER_OF(2) entries: Shifted by N.
-// (***) ATTACH entry: No MEMBER_OF bit; no inherited map-type-modifying bits
-// (DELETE/ALWAYS/CLOSE/PRESENT).
-// (i = array element index; N = __tgt_mapper_num_components()):
+// (*) Pointee (combined entry): No MEMBER_OF addition; modifiers (ALWAYS/DELETE/CLOSE/PRESENT)
+// propagated from the outer map type.
+// (**) Inner MEMBER_OF(2) entries: Shifted by N; modifiers (ALWAYS/DELETE/CLOSE/PRESENT)
+// propagated from the outer map type.
+// (***) ATTACH entry: No MEMBER_OF bit; no propagated modifiers.
typedef struct {
int x;
diff --git a/mlir/test/Target/LLVMIR/omptarget-llvm.mlir b/mlir/test/Target/LLVMIR/omptarget-llvm.mlir
index 6990ea001b6e4..ee8acf0037eba 100644
--- a/mlir/test/Target/LLVMIR/omptarget-llvm.mlir
+++ b/mlir/test/Target/LLVMIR/omptarget-llvm.mlir
@@ -603,7 +603,9 @@ module attributes {omp.target_triples = ["amdgcn-amd-amdhsa"]} {
// CHECK: br label %[[VAL_42]]
// CHECK: omp.type.end: ; preds = %[[VAL_59]], %[[VAL_56]], %[[VAL_55]], %[[VAL_51]]
// CHECK: %[[VAL_61:.*]] = phi i64 [ %[[VAL_53]], %[[VAL_51]] ], [ %[[VAL_57]], %[[VAL_55]] ], [ %[[VAL_60]], %[[VAL_59]] ], [ %[[VAL_48]], %[[VAL_56]] ]
-// CHECK: call void @__tgt_push_mapper_component(ptr %[[VAL_37]], ptr %[[VAL_43]], ptr %[[VAL_45]], i64 4, i64 %[[VAL_61]], ptr @2)
+// CHECK: %[[VAL_MODMASK:.*]] = and i64 %[[VAL_22]], 5132
+// CHECK: %[[VAL_WITH_MOD:.*]] = or i64 %[[VAL_61]], %[[VAL_MODMASK]]
+// CHECK: call void @__tgt_push_mapper_component(ptr %[[VAL_37]], ptr %[[VAL_43]], ptr %[[VAL_45]], i64 4, i64 %[[VAL_WITH_MOD]], ptr @2)
// CHECK: %[[VAL_44]] = getelementptr %[[VAL_18]], ptr %[[VAL_43]], i32 1
// CHECK: %[[VAL_62:.*]] = icmp eq ptr %[[VAL_44]], %[[VAL_17]]
// CHECK: br i1 %[[VAL_62]], label %[[VAL_63:.*]], label %[[VAL_41]]
diff --git a/offload/test/mapping/mapper_map_mbr_then_present_mbr_ptee.c b/offload/test/mapping/mapper_map_mbr_then_present_mbr_ptee.c
index 0939aa0b2330f..8550585495511 100644
--- a/offload/test/mapping/mapper_map_mbr_then_present_mbr_ptee.c
+++ b/offload/test/mapping/mapper_map_mbr_then_present_mbr_ptee.c
@@ -5,7 +5,7 @@
#include <stdio.h>
// Check that the present check fails if we map a struct member, then do a
-// do a map(present) on a mapper that maps both a member and a pointee.
+// map(present) on a mapper that maps both a member and a pointee.
int x[10];
diff --git a/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections.c b/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections.c
index 112f60165361f..c9c1a2c8ae7a9 100644
--- a/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections.c
+++ b/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections.c
@@ -3,8 +3,9 @@
#include <omp.h>
#include <stdio.h>
-// Check that the mapping of members of a struct and pointee data using a mapper
-// properly allocates/deletes all fields.
+// The mapper maps s2.z, s2.s1p->x, s2.s1p->y, and s2.s1p->p[0:10].
+// Check that s2.s1p->dummy and s2.s1p->p itself are not mapped, and that all
+// mapped fields are correctly removed on delete.
int x[10];
diff --git a/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections_array.c b/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections_array.c
index 076867f05470e..2f1c5a9a7e615 100644
--- a/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections_array.c
+++ b/offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections_array.c
@@ -4,9 +4,12 @@
#include <stdio.h>
// Array variant of mapper_map_ptee_only_2_ptr_indirections.c.
-// Exercises the nested-pointer-chain case (s2.s1p->p[0:10]) in an array
-// context, where inner MEMBER_OF bits must be shifted correctly and outer
-// MEMBER_OF must not be applied to the pointee entries.
+// The mapper maps s2.z, s2.s1p->x, s2.s1p->y, and s2.s1p->p[0:10].
+// s2.s1p->dummy and s2.s1p->p itself are not mapped.
+// This exercises the nested-pointer-chain case: the inner MEMBER_OF bits for
+// s2.s1p->x/y/p[0:10] must be shifted correctly, and outer MEMBER_OF must not
+// be applied to the pointee entry (s2.s1p->p[0:10]) or the ATTACH entry
+// (s2.s1p).
int x[2][10];
diff --git a/offload/test/mapping/mapper_map_ptee_only_2ndlevel_array.c b/offload/test/mapping/mapper_map_ptee_only_2ndlevel_array.c
index 08b9551f0c8a9..c745779ab4b97 100644
--- a/offload/test/mapping/mapper_map_ptee_only_2ndlevel_array.c
+++ b/offload/test/mapping/mapper_map_ptee_only_2ndlevel_array.c
@@ -3,9 +3,13 @@
#include <omp.h>
#include <stdio.h>
-// Test that a mapper on a nested struct maps the right members when applied to
-// an array of structs: s.x and s.p[0:10] are mapped; s.dummy and s.p itself
-// are not (modulo attach FIXME).
+// Test that a mapper on a nested struct correctly maps and deletes data when
+// applied to an array of structs. s2arr[i].s1.x and s2arr[i].s1.p[0:10] are
+// mapped; after delete they are gone.
+// It's not ideal that s2arr[i].s1.dummy and s2arr[i].s1.p are also present
+// after enter-data (the runtime allocates the full contiguous storage of the
+// array, otherwise it's tricky to map an array with "holes"), but everything is
+// correctly removed after delete.
int x[2][10];
>From 94c8011e4e69798809cb608ad95d46bd8f15825d Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Wed, 13 May 2026 18:13:48 -0700
Subject: [PATCH 12/13] Minor comment change.
---
.../target_map_array_of_structs_with_nested_mapper_codegen.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_codegen.cpp b/clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_codegen.cpp
index 70514efe174af..b14d19ac83187 100644
--- a/clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_codegen.cpp
+++ b/clang/test/OpenMP/target_map_array_of_structs_with_nested_mapper_codegen.cpp
@@ -273,7 +273,7 @@ void foo() {
//
// C mapper for: map(to: s.a) — per-element entries for struct C
// (i = array element index, N = __tgt_mapper_num_components()):
-// &c[i], &c[i].a, sizeof(int), MEMBER_OF(N) | TO | FROM | modifiers
+// &c[i], &c[i].a, sizeof(int), MEMBER_OF(N) | TO | modifiers
//
// CHECK-LABEL: define {{[^@]+}}@.omp_mapper._ZTS1C.default
// CHECK-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]], ptr noundef [[TMP2:%.*]], i64 noundef [[TMP3:%.*]], i64 noundef [[TMP4:%.*]], ptr noundef [[TMP5:%.*]]) #[[ATTR2]] {
>From 88fc2a0fc2b760d9620d44392b77f282a9eb15fa Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Thu, 14 May 2026 18:25:48 -0700
Subject: [PATCH 13/13] Update the new struct from mlir translator.
---
.../LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp | 7 +++++++
1 file changed, 7 insertions(+)
diff --git a/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp b/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp
index 5bb97d8fd6588..54323b5ea9e93 100644
--- a/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp
+++ b/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp
@@ -5556,6 +5556,7 @@ static llvm::omp::OpenMPOffloadMappingFlags mapParentWithMembers(
}
combinedInfo.Types.emplace_back(baseFlag);
+ combinedInfo.DontAddMemberOfInMapper.emplace_back(false);
combinedInfo.DevicePointers.emplace_back(
mapData.DevicePointers[mapDataIndex]);
// Only attach the mapper to the base entry when we are mapping the whole
@@ -5626,6 +5627,7 @@ static llvm::omp::OpenMPOffloadMappingFlags mapParentWithMembers(
if (targetDirective == TargetDirectiveEnumTy::TargetUpdate || hasMapClose) {
combinedInfo.Types.emplace_back(mapFlag);
+ combinedInfo.DontAddMemberOfInMapper.emplace_back(false);
combinedInfo.DevicePointers.emplace_back(
mapData.DevicePointers[mapDataIndex]);
combinedInfo.Names.emplace_back(LLVM::createMappingInformation(
@@ -5662,6 +5664,7 @@ static llvm::omp::OpenMPOffloadMappingFlags mapParentWithMembers(
mapData,
cast<omp::MapInfoOp>(parentClause.getMembers()[v].getDefiningOp()));
combinedInfo.Types.emplace_back(mapFlag);
+ combinedInfo.DontAddMemberOfInMapper.emplace_back(false);
combinedInfo.DevicePointers.emplace_back(
mapData.DevicePointers[mapDataOverlapIdx]);
combinedInfo.Names.emplace_back(LLVM::createMappingInformation(
@@ -5684,6 +5687,7 @@ static llvm::omp::OpenMPOffloadMappingFlags mapParentWithMembers(
}
combinedInfo.Types.emplace_back(mapFlag);
+ combinedInfo.DontAddMemberOfInMapper.emplace_back(false);
combinedInfo.DevicePointers.emplace_back(
mapData.DevicePointers[mapDataIndex]);
combinedInfo.Names.emplace_back(LLVM::createMappingInformation(
@@ -5733,6 +5737,7 @@ static void processMapMembersWithParent(
mapFlag |= llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_MEMBER_OF;
ompBuilder.setCorrectMemberOfFlag(mapFlag, memberOfFlag);
combinedInfo.Types.emplace_back(mapFlag);
+ combinedInfo.DontAddMemberOfInMapper.emplace_back(false);
combinedInfo.DevicePointers.emplace_back(
llvm::OpenMPIRBuilder::DeviceInfoTy::None);
combinedInfo.Mappers.emplace_back(nullptr);
@@ -5762,6 +5767,7 @@ static void processMapMembersWithParent(
}
combinedInfo.Types.emplace_back(mapFlag);
+ combinedInfo.DontAddMemberOfInMapper.emplace_back(false);
combinedInfo.DevicePointers.emplace_back(
mapData.DevicePointers[memberDataIdx]);
combinedInfo.Mappers.emplace_back(mapData.Mappers[memberDataIdx]);
@@ -5820,6 +5826,7 @@ static void processIndividualMap(MapInfoData &mapData, size_t mapDataIdx,
combinedInfo.Mappers.emplace_back(mapData.Mappers[mapDataIdx]);
combinedInfo.Names.emplace_back(mapData.Names[mapDataIdx]);
combinedInfo.Types.emplace_back(mapFlag);
+ combinedInfo.DontAddMemberOfInMapper.emplace_back(false);
combinedInfo.Sizes.emplace_back(mapData.Sizes[mapDataIdx]);
}
More information about the Mlir-commits
mailing list