[clang] [llvm] [mlir] [OpenMP][Clang] Enable `ATTACH`-style maps for mappers. (PR #210213)
Abhinav Gaba via llvm-commits
llvm-commits at lists.llvm.org
Sun Aug 2 23:25:53 PDT 2026
https://github.com/abhinavgaba updated https://github.com/llvm/llvm-project/pull/210213
>From dc7adf8f5e7dc746e280cc00f2ecc121ab3fbe4e Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Mon, 13 Jul 2026 20:31:45 -0700
Subject: [PATCH 1/2] [OpenMP][Clang] Enable ATTACH-style maps for mappers.
Track per-entry attach-ptr info (HasAttachPtr) through mapper codegen so that
emitUserDefinedMapper does not add a new outer MEMBER_OF to pointee/combined
entries (which occupy different storage than the struct) or to ATTACH entries.
Clang and the MLIR translator populate the per-entry array in parallel with the
other MapInfosTy arrays.
Address review:
- Rename MapSkipMemberOfArrayTy to MapHasAttachPtrArrayTy to match the
HasAttachPtr field it backs.
- Restructure the emitUserDefinedMapper comment into a bulleted (*)/(**)/(***)
list keyed to the example entries.
- Reword the Clang comments: HasAttachPtr marks pointee entries that have a
base attach-ptr; a combined entry has a base attach-ptr if its constituents
do; cross-reference emitUserDefinedMapper for the MEMBER_OF rationale.
- Update the moved present-check tests to their now-correct behavior (the
attach-style maps make the inbounds present checks pass and remove the
"explicit extension" errors).
Co-Authored-By: Claude Opus 4.8 <noreply at anthropic.com>
---
clang/docs/OpenMPSupport.md | 5 +-
clang/docs/ReleaseNotes.md | 3 +
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 70 +++-
clang/test/OpenMP/declare_mapper_codegen.cpp | 128 ++++---
...t_map_nested_ptr_member_mapper_codegen.cpp | 356 ++++++++++--------
.../llvm/Frontend/OpenMP/OMPIRBuilder.h | 6 +
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 71 +++-
.../Frontend/OpenMPIRBuilderTest.cpp | 1 +
.../OpenMP/OpenMPToLLVMIRTranslation.cpp | 20 +
.../mapper_enter_data_always_present_ptee.c | 42 +--
...apper_map_mbr_ptee_then_present_mbr_ptee.c | 42 +--
.../mapper_map_mbr_then_present_mbr_ptee.c | 40 +-
.../test/mapping/mapper_map_present_ptee.c | 26 +-
offload/test/mapping/mapper_map_ptee_only.c | 14 +-
.../mapper_map_ptee_only_2_ptr_indirections.c | 6 +-
...r_map_ptee_only_2_ptr_indirections_array.c | 14 +-
.../mapping/mapper_map_ptee_only_2ndlevel.c | 12 +-
17 files changed, 494 insertions(+), 362 deletions(-)
diff --git a/clang/docs/OpenMPSupport.md b/clang/docs/OpenMPSupport.md
index ff9dd01978040..840cc5ef0fecd 100644
--- a/clang/docs/OpenMPSupport.md
+++ b/clang/docs/OpenMPSupport.md
@@ -134,7 +134,7 @@ implementation.
| device | support close modifier on map clause | {good}`done` | [D55719][D55719],[D55892][D55892] |
| device | teams construct on the host device | {good}`done` | r371553 |
| device | support non-contiguous array sections for target update | {good}`done` | [PR144635][PR144635] |
-| device | pointer attachment | {part}`being repaired` | @abhinavgaba ([PR153683][PR153683]) |
+| device | pointer attachment | {part}`being repaired` | @abhinavgaba ([PR153683][PR153683], [PR210213][PR210213]) |
| atomic | hints for the atomic construct | {good}`done` | [D51233][D51233] |
| base language | C11 support | {good}`done` | |
| base language | C++11/14/17 support | {good}`done` | |
@@ -386,7 +386,7 @@ implementation.
| dyn_groupprivate clause | {part}`partial` | {part}`In Progress` | C/C++: Host device support missing |
| loop flatten transformation | {none}`unclaimed` | {none}`unclaimed` | |
| loop grid/tile modifiers for sizes clause | {none}`unclaimed` | {none}`unclaimed` | |
-| attach map-type modifier | {part}`In Progress` | {none}`unclaimed` | C/C++: @abhinavgaba; RT: @abhinavgaba ([PR149036][PR149036], [PR158370][PR158370]) |
+| attach map-type modifier | {part}`In Progress` | {none}`unclaimed` | C/C++: @abhinavgaba; RT: @abhinavgaba ([PR149036][PR149036], [PR158370][PR158370], [PR210213][PR210213]) |
| need_device_ptr modifier for adjust_args clause | {part}`partial` | {none}`unclaimed` | Clang Parsing/Sema: [PR168905][PR168905] [PR169558][PR169558] |
| fallback modifier for use_device_ptr clause | {good}`done` | {none}`unclaimed` | Clang: @abhinavgaba ([PR170578][PR170578], [PR173931][PR173931]) RT: @abhinavgaba ([PR169603][PR169603]) |
| dims modifier for num_teams, thread_limit, and num_threads clauses | {part}`partial` | {part}`In Progress` | C/C++: @kevinsala ([PR206412]); Fortran: @skc7, @kparzysz, @mjklemm |
@@ -562,5 +562,6 @@ considered for standardization. Please post on the
[PR194168]: https://github.com/llvm/llvm-project/pull/194168
[PR195829]: https://github.com/llvm/llvm-project/pull/195829
[PR196431]: https://github.com/llvm/llvm-project/pull/196431
+[PR210213]: https://github.com/llvm/llvm-project/pull/210213
[discourse forums (runtimes - openmp category)]: https://discourse.llvm.org/c/runtimes/openmp/35
diff --git a/clang/docs/ReleaseNotes.md b/clang/docs/ReleaseNotes.md
index 7108392abbaa1..321bd7cd24ac0 100644
--- a/clang/docs/ReleaseNotes.md
+++ b/clang/docs/ReleaseNotes.md
@@ -548,6 +548,9 @@ features cannot lower the translation-unit ABI level;
`thread_limit` clauses for OpenMP 6.1 or later.
- Map-type-modifying modifiers applied to a list item with a user-defined mapper
are now propagated onto the maps the mapper expands to.
+- Mapping of expressions with base-pointers through a user-defined mapper (e.g.
+ `map(s.p[0:n])`) now conforms to OpenMP's conditional pointer-attachment,
+ matching the behavior of such maps outside a mapper.
### SYCL Support
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 9e71e82e5ee86..6edabd4e42d88 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7605,6 +7605,8 @@ class MappableExprsHandler {
AttachInfo.AttachPteeAddr.emitRawPointer(CGF));
CombinedInfo.Sizes.push_back(PointerSize);
CombinedInfo.Types.push_back(OpenMPOffloadMappingFlags::OMP_MAP_ATTACH);
+ // ATTACH entries themselves don't "have" a base attach-ptr.
+ CombinedInfo.HasAttachPtr.push_back(false);
CombinedInfo.Mappers.push_back(nullptr);
CombinedInfo.NonContigInfo.Dims.push_back(1);
}
@@ -7706,6 +7708,7 @@ class MappableExprsHandler {
CombinedInfo.Sizes.push_back(
CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/false));
CombinedInfo.Types.push_back(Flags);
+ CombinedInfo.HasAttachPtr.push_back(false);
CombinedInfo.Mappers.push_back(nullptr);
CombinedInfo.NonContigInfo.Dims.push_back(IsNonContiguous ? DimSize : 1);
}
@@ -8393,10 +8396,14 @@ class MappableExprsHandler {
}
}
- if (!IsMappingWholeStruct)
+ if (!IsMappingWholeStruct) {
CombinedInfo.Types.push_back(Flags);
- else
+ // HasAttachPtr marks pointee entries, which have a base attach-ptr.
+ CombinedInfo.HasAttachPtr.push_back(HasAttachPtr);
+ } else {
StructBaseCombinedInfo.Types.push_back(Flags);
+ StructBaseCombinedInfo.HasAttachPtr.push_back(HasAttachPtr);
+ }
}
// If we have encountered a member expression so far, keep track of the
@@ -8997,6 +9004,7 @@ class MappableExprsHandler {
if (HasUdpFbNullify)
Flags |= OpenMPOffloadMappingFlags::OMP_MAP_FB_NULLIFY;
UseDeviceDataCombinedInfo.Types.push_back(Flags);
+ UseDeviceDataCombinedInfo.HasAttachPtr.push_back(false);
UseDeviceDataCombinedInfo.Mappers.push_back(nullptr);
};
@@ -9404,7 +9412,28 @@ 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, to enable attach-style mapping
+ // for mappers.
+ 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
@@ -9478,6 +9507,14 @@ class MappableExprsHandler {
: !PartialStruct.PreliminaryMapData.BasePointers.empty()
? OpenMPOffloadMappingFlags::OMP_MAP_PTR_AND_OBJ
: OpenMPOffloadMappingFlags::OMP_MAP_TARGET_PARAM);
+ // A combined entry has a base attach-ptr if its constituents do. e.g.:
+ // map(s2.s1p->x, s2.s1p->y)
+ // combined entry:
+ // s2.s1p[0], s2.s1p->x, sizeof(s1p->x..y), ALLOC
+ // here s2.s1p is the attach-ptr for the combined entry.
+ // See the inline comments in emitUserDefinedMapper's definition for how
+ // entries with an attach-ptr are treated.
+ CombinedInfo.HasAttachPtr.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() !=
@@ -9593,6 +9630,7 @@ class MappableExprsHandler {
OpenMPOffloadMappingFlags::OMP_MAP_LITERAL |
OpenMPOffloadMappingFlags::OMP_MAP_MEMBER_OF |
OpenMPOffloadMappingFlags::OMP_MAP_IMPLICIT);
+ CombinedInfo.HasAttachPtr.push_back(false);
CombinedInfo.Mappers.push_back(nullptr);
}
for (const LambdaCapture &LC : RD->captures()) {
@@ -9633,6 +9671,7 @@ class MappableExprsHandler {
OpenMPOffloadMappingFlags::OMP_MAP_LITERAL |
OpenMPOffloadMappingFlags::OMP_MAP_MEMBER_OF |
OpenMPOffloadMappingFlags::OMP_MAP_IMPLICIT);
+ CombinedInfo.HasAttachPtr.push_back(false);
CombinedInfo.Mappers.push_back(nullptr);
}
}
@@ -9925,6 +9964,7 @@ class MappableExprsHandler {
CurCaptureVarInfo.Types.push_back(
OpenMPOffloadMappingFlags::OMP_MAP_LITERAL |
OpenMPOffloadMappingFlags::OMP_MAP_TARGET_PARAM);
+ CurCaptureVarInfo.HasAttachPtr.push_back(false);
CurCaptureVarInfo.Mappers.push_back(nullptr);
return;
}
@@ -10323,6 +10363,7 @@ class MappableExprsHandler {
if (IsImplicit)
CombinedInfo.Types.back() |= OpenMPOffloadMappingFlags::OMP_MAP_IMPLICIT;
+ CombinedInfo.HasAttachPtr.push_back(false);
// No user-defined mapper for default mapping.
CombinedInfo.Mappers.push_back(nullptr);
}
@@ -10539,14 +10580,31 @@ 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) 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). PRESENT is handled separately.
+/// imported_modifier_bits = type & (ALWAYS | DELETE | CLOSE);
+/// effective_type = c.isAttach() ? member_type
+/// : member_type | imported_modifier_bits;
/// 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);
/// }
/// }
@@ -10762,6 +10820,7 @@ static void genMapInfoForCaptures(
CurInfo.Types.push_back(OpenMPOffloadMappingFlags::OMP_MAP_LITERAL |
OpenMPOffloadMappingFlags::OMP_MAP_TARGET_PARAM |
OpenMPOffloadMappingFlags::OMP_MAP_IMPLICIT);
+ CurInfo.HasAttachPtr.push_back(false);
CurInfo.Mappers.push_back(nullptr);
} else {
const ValueDecl *CapturedVD =
@@ -10919,6 +10978,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.HasAttachPtr.push_back(false);
if (!CombinedInfo.Names.empty())
CombinedInfo.Names.push_back(NullPtr);
CombinedInfo.Exprs.push_back(nullptr);
diff --git a/clang/test/OpenMP/declare_mapper_codegen.cpp b/clang/test/OpenMP/declare_mapper_codegen.cpp
index eb90f218f5ca8..3eb4c2f28efb2 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:[^,]+]]
@@ -150,57 +149,48 @@ class C {
// CK0-DAG: [[PHITYPE0:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
// CK0-DAG: [[MODMASK0:%.+]] = and i64 [[TYPE]], 1036
// CK0-DAG: [[PHITYPE0_MOD:%.+]] = or i64 [[PHITYPE0]], [[MODMASK0]]
-// CK0: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 [[CUSIZE]], i64 [[PHITYPE0_MOD]], {{.*}})
-// 281474976710659 == 0x1,000,000,003
-// CK0-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710659, [[SHIPRESIZE]]
+// 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-DAG: [[TYPE1:%.+]] = phi i64 [ 0, %[[ALLOC]] ], [ 1, %[[TO]] ], [ 2, %[[FROM]] ], [ 3, %[[TOELSE]] ]
// CK0-DAG: [[MODMASK1:%.+]] = and i64 [[TYPE]], 1036
// CK0-DAG: [[TYPE1_MOD:%.+]] = or i64 [[TYPE1]], [[MODMASK1]]
-// CK0: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 4, i64 [[TYPE1_MOD]], {{.*}})
-// 281474976710675 == 0x1,000,000,013
-// CK0-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710675, [[SHIPRESIZE]]
+// 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-DAG: [[MODMASK2:%.+]] = and i64 [[TYPE]], 1036
-// CK0-DAG: [[TYPE2_MOD:%.+]] = or i64 [[TYPE2]], [[MODMASK2]]
-// CK0: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BBEGIN]], ptr [[BARRBEGINGEP]], i64 16, i64 [[TYPE2_MOD]], {{.*}})
+// 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]]
@@ -592,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
@@ -699,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{{.*}})
@@ -894,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
@@ -924,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:[^,]+]]
@@ -960,57 +955,48 @@ class C {
// CK4-DAG: [[PHITYPE0:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
// CK4-DAG: [[MODMASK0:%.+]] = and i64 [[TYPE]], 1036
// CK4-DAG: [[PHITYPE0_MOD:%.+]] = or i64 [[PHITYPE0]], [[MODMASK0]]
-// CK4: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 [[CUSIZE]], i64 [[PHITYPE0_MOD]], {{.*}})
-// 281474976710659 == 0x1,000,000,003
-// CK4-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710659, [[SHIPRESIZE]]
+// 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-DAG: [[TYPE1:%.+]] = phi i64 [ 0, %[[ALLOC]] ], [ 1, %[[TO]] ], [ 2, %[[FROM]] ], [ 3, %[[TOELSE]] ]
// CK4-DAG: [[MODMASK1:%.+]] = and i64 [[TYPE]], 1036
// CK4-DAG: [[TYPE1_MOD:%.+]] = or i64 [[TYPE1]], [[MODMASK1]]
-// CK4: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 4, i64 [[TYPE1_MOD]], {{.*}})
-// 281474976710675 == 0x1,000,000,013
-// CK4-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710675, [[SHIPRESIZE]]
+// 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-DAG: [[MODMASK2:%.+]] = and i64 [[TYPE]], 1036
-// CK4-DAG: [[TYPE2_MOD:%.+]] = or i64 [[TYPE2]], [[MODMASK2]]
-// CK4: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BBEGIN]], ptr [[BARRBEGINGEP]], i64 16, i64 [[TYPE2_MOD]], {{.*}})
+// 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]]
@@ -1083,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)
{
@@ -1117,34 +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: [[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-DAG: [[TYPE1:%.+]] = phi i64 [ 0, %[[ALLOC]] ], [ 1, %[[TO]] ], [ 2, %[[FROM]] ], [ 3, %[[TOELSE]] ]
// CK5-DAG: [[MODMASK:%.+]] = and i64 [[TYPE]], 1036
// CK5-DAG: [[TYPE1_MOD:%.+]] = or i64 [[TYPE1]], [[MODMASK]]
-// CK5: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 {{.*}}, i64 [[TYPE1_MOD]], {{.*}})
+// 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: 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: [[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_nested_ptr_member_mapper_codegen.cpp b/clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp
index bd9189dce7c9f..821b5fd1652c4 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
@@ -98,132 +98,150 @@ void foo(S2 *arr) {
// 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 [ [[TMP2]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END25:%.*]] ]
+// 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: [[TMP15:%.*]] = load ptr, ptr [[S1P1]], align 8
-// CHECK: [[X:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP15]], 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: [[TMP16:%.*]] = load ptr, ptr [[S1P3]], align 8
-// CHECK: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S1]], ptr [[TMP16]], i32 0, i32 1
-// CHECK: [[TMP17:%.*]] = getelementptr i32, ptr [[Z]], i32 1
-// CHECK: [[TMP18:%.*]] = ptrtoaddr ptr [[TMP17]] to i64
-// CHECK: [[TMP19:%.*]] = ptrtoaddr ptr [[S1P]] to i64
-// CHECK: [[TMP20:%.*]] = sub i64 [[TMP18]], [[TMP19]]
-// CHECK: [[TMP21:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP0]])
-// CHECK: [[TMP22:%.*]] = shl i64 [[TMP21]], 48
-// CHECK: [[TMP23:%.*]] = add nuw i64 0, [[TMP22]]
-// CHECK: [[TMP24:%.*]] = and i64 [[TMP4]], 3
-// CHECK: [[TMP25:%.*]] = icmp eq i64 [[TMP24]], 0
-// CHECK: br i1 [[TMP25]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]]
+// 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: [[TMP26:%.*]] = and i64 [[TMP23]], -4
+// CHECK: [[TMP28:%.*]] = and i64 [[TMP25]], -4
// CHECK: br label [[OMP_TYPE_END:%.*]]
// CHECK: omp.type.alloc.else:
-// CHECK: [[TMP27:%.*]] = icmp eq i64 [[TMP24]], 1
-// CHECK: br i1 [[TMP27]], 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: [[TMP28:%.*]] = and i64 [[TMP23]], -3
+// CHECK: [[TMP30:%.*]] = and i64 [[TMP25]], -3
// CHECK: br label [[OMP_TYPE_END]]
// CHECK: omp.type.to.else:
-// CHECK: [[TMP29:%.*]] = icmp eq i64 [[TMP24]], 2
-// CHECK: br i1 [[TMP29]], 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: [[TMP30:%.*]] = and i64 [[TMP23]], -2
+// CHECK: [[TMP32:%.*]] = and i64 [[TMP25]], -2
// CHECK: br label [[OMP_TYPE_END]]
// CHECK: omp.type.end:
-// CHECK: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP26]], [[OMP_TYPE_ALLOC]] ], [ [[TMP28]], [[OMP_TYPE_TO]] ], [ [[TMP30]], [[OMP_TYPE_FROM]] ], [ [[TMP23]], [[OMP_TYPE_TO_ELSE]] ]
-// CHECK: [[TMP31:%.*]] = and i64 [[TMP4]], 1036
-// CHECK: [[OMP_MAPTYPE_WITH_MODIFIERS:%.*]] = or i64 [[OMP_MAPTYPE]], [[TMP31]]
-// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[S1P]], i64 [[TMP20]], i64 [[OMP_MAPTYPE_WITH_MODIFIERS]], ptr null)
-// CHECK: [[TMP32:%.*]] = add nuw i64 281474976710659, [[TMP22]]
-// CHECK: [[TMP33:%.*]] = and i64 [[TMP4]], 3
-// CHECK: [[TMP34:%.*]] = icmp eq i64 [[TMP33]], 0
-// CHECK: br i1 [[TMP34]], label [[OMP_TYPE_ALLOC4:%.*]], label [[OMP_TYPE_ALLOC_ELSE5:%.*]]
+// 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]], 1036
+// 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: [[TMP35:%.*]] = and i64 [[TMP32]], -4
// CHECK: br label [[OMP_TYPE_END9:%.*]]
// CHECK: omp.type.alloc.else5:
-// CHECK: [[TMP36:%.*]] = icmp eq i64 [[TMP33]], 1
+// 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: [[TMP37:%.*]] = and i64 [[TMP32]], -3
// CHECK: br label [[OMP_TYPE_END9]]
// CHECK: omp.type.to.else7:
-// CHECK: [[TMP38:%.*]] = icmp eq i64 [[TMP33]], 2
-// CHECK: br i1 [[TMP38]], label [[OMP_TYPE_FROM8:%.*]], label [[OMP_TYPE_END9]]
+// CHECK: [[TMP37:%.*]] = icmp eq i64 [[TMP34]], 2
+// CHECK: br i1 [[TMP37]], label [[OMP_TYPE_FROM8:%.*]], label [[OMP_TYPE_END9]]
// CHECK: omp.type.from8:
-// CHECK: [[TMP39:%.*]] = and i64 [[TMP32]], -2
// CHECK: br label [[OMP_TYPE_END9]]
// CHECK: omp.type.end9:
-// CHECK: [[OMP_MAPTYPE10:%.*]] = phi i64 [ [[TMP35]], [[OMP_TYPE_ALLOC4]] ], [ [[TMP37]], [[OMP_TYPE_TO6]] ], [ [[TMP39]], [[OMP_TYPE_FROM8]] ], [ [[TMP32]], [[OMP_TYPE_TO_ELSE7]] ]
-// CHECK: [[TMP40:%.*]] = and i64 [[TMP4]], 1036
-// CHECK: [[OMP_MAPTYPE_WITH_MODIFIERS11:%.*]] = or i64 [[OMP_MAPTYPE10]], [[TMP40]]
-// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[Z]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS11]], ptr null)
-// CHECK: [[TMP41:%.*]] = add nuw i64 281474976710675, [[TMP22]]
-// CHECK: [[TMP42:%.*]] = and i64 [[TMP4]], 3
-// CHECK: [[TMP43:%.*]] = icmp eq i64 [[TMP42]], 0
-// CHECK: br i1 [[TMP43]], label [[OMP_TYPE_ALLOC12:%.*]], label [[OMP_TYPE_ALLOC_ELSE13:%.*]]
+// 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]], 1036
+// 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: [[TMP44:%.*]] = and i64 [[TMP41]], -4
+// CHECK: [[TMP42:%.*]] = and i64 [[TMP39]], -4
// CHECK: br label [[OMP_TYPE_END17:%.*]]
// CHECK: omp.type.alloc.else13:
-// CHECK: [[TMP45:%.*]] = icmp eq i64 [[TMP42]], 1
-// CHECK: br i1 [[TMP45]], label [[OMP_TYPE_TO14:%.*]], label [[OMP_TYPE_TO_ELSE15:%.*]]
+// 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: [[TMP46:%.*]] = and i64 [[TMP41]], -3
+// CHECK: [[TMP44:%.*]] = and i64 [[TMP39]], -3
// CHECK: br label [[OMP_TYPE_END17]]
// CHECK: omp.type.to.else15:
-// CHECK: [[TMP47:%.*]] = icmp eq i64 [[TMP42]], 2
-// CHECK: br i1 [[TMP47]], label [[OMP_TYPE_FROM16:%.*]], label [[OMP_TYPE_END17]]
+// CHECK: [[TMP45:%.*]] = icmp eq i64 [[TMP40]], 2
+// CHECK: br i1 [[TMP45]], label [[OMP_TYPE_FROM16:%.*]], label [[OMP_TYPE_END17]]
// CHECK: omp.type.from16:
-// CHECK: [[TMP48:%.*]] = and i64 [[TMP41]], -2
+// CHECK: [[TMP46:%.*]] = and i64 [[TMP39]], -2
// CHECK: br label [[OMP_TYPE_END17]]
// CHECK: omp.type.end17:
-// CHECK: [[OMP_MAPTYPE18:%.*]] = phi i64 [ [[TMP44]], [[OMP_TYPE_ALLOC12]] ], [ [[TMP46]], [[OMP_TYPE_TO14]] ], [ [[TMP48]], [[OMP_TYPE_FROM16]] ], [ [[TMP41]], [[OMP_TYPE_TO_ELSE15]] ]
-// CHECK: [[TMP49:%.*]] = and i64 [[TMP4]], 1036
-// CHECK: [[OMP_MAPTYPE_WITH_MODIFIERS19:%.*]] = or i64 [[OMP_MAPTYPE18]], [[TMP49]]
-// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[S1P]], ptr [[X]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS19]], ptr null)
-// CHECK: [[TMP50:%.*]] = add nuw i64 281474976710675, [[TMP22]]
-// CHECK: [[TMP51:%.*]] = and i64 [[TMP4]], 3
-// CHECK: [[TMP52:%.*]] = icmp eq i64 [[TMP51]], 0
-// CHECK: br i1 [[TMP52]], label [[OMP_TYPE_ALLOC20:%.*]], label [[OMP_TYPE_ALLOC_ELSE21:%.*]]
+// 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]], 1036
+// 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: [[TMP53:%.*]] = and i64 [[TMP50]], -4
-// CHECK: br label [[OMP_TYPE_END25]]
+// CHECK: [[TMP51:%.*]] = and i64 [[TMP48]], -4
+// CHECK: br label [[OMP_TYPE_END25:%.*]]
// CHECK: omp.type.alloc.else21:
-// CHECK: [[TMP54:%.*]] = icmp eq i64 [[TMP51]], 1
-// CHECK: br i1 [[TMP54]], label [[OMP_TYPE_TO22:%.*]], label [[OMP_TYPE_TO_ELSE23:%.*]]
+// 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: [[TMP55:%.*]] = and i64 [[TMP50]], -3
+// CHECK: [[TMP53:%.*]] = and i64 [[TMP48]], -3
// CHECK: br label [[OMP_TYPE_END25]]
// CHECK: omp.type.to.else23:
-// CHECK: [[TMP56:%.*]] = icmp eq i64 [[TMP51]], 2
-// CHECK: br i1 [[TMP56]], label [[OMP_TYPE_FROM24:%.*]], label [[OMP_TYPE_END25]]
+// CHECK: [[TMP54:%.*]] = icmp eq i64 [[TMP49]], 2
+// CHECK: br i1 [[TMP54]], label [[OMP_TYPE_FROM24:%.*]], label [[OMP_TYPE_END25]]
// CHECK: omp.type.from24:
-// CHECK: [[TMP57:%.*]] = and i64 [[TMP50]], -2
+// CHECK: [[TMP55:%.*]] = and i64 [[TMP48]], -2
// CHECK: br label [[OMP_TYPE_END25]]
// CHECK: omp.type.end25:
-// CHECK: [[OMP_MAPTYPE26:%.*]] = phi i64 [ [[TMP53]], [[OMP_TYPE_ALLOC20]] ], [ [[TMP55]], [[OMP_TYPE_TO22]] ], [ [[TMP57]], [[OMP_TYPE_FROM24]] ], [ [[TMP50]], [[OMP_TYPE_TO_ELSE23]] ]
-// CHECK: [[TMP58:%.*]] = and i64 [[TMP4]], 1036
-// CHECK: [[OMP_MAPTYPE_WITH_MODIFIERS27:%.*]] = or i64 [[OMP_MAPTYPE26]], [[TMP58]]
-// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[S1P2]], ptr [[Y]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS27]], ptr null)
+// 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]], 1036
+// 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]], 1036
+// 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]], [[TMP7]]
// CHECK: br i1 [[OMP_ARRAYMAP_ISDONE]], label [[OMP_ARRAYMAP_EXIT:%.*]], label [[OMP_ARRAYMAP_BODY]]
// CHECK: omp.arraymap.exit:
-// CHECK: [[OMP_ARRAYINIT_ISARRAY28:%.*]] = icmp sgt i64 [[TMP6]], 1
-// CHECK: [[TMP59:%.*]] = and i64 [[TMP4]], 8
-// CHECK: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP59]], 0
-// CHECK: [[TMP60:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY28]], [[DOTOMP_ARRAY__DEL__DELETE]]
-// CHECK: br i1 [[TMP60]], 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: [[TMP61:%.*]] = mul nuw i64 [[TMP6]], 16
-// CHECK: [[TMP62:%.*]] = and i64 [[TMP4]], -4
-// CHECK: [[TMP63:%.*]] = or i64 [[TMP62]], 512
-// CHECK: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP61]], i64 [[TMP63]], ptr [[TMP5]])
+// 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
@@ -276,132 +294,150 @@ void foo(S2 *arr) {
// CHECK-60: [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP2]], [[TMP7]]
// CHECK-60: br i1 [[OMP_ARRAYMAP_ISEMPTY]], label [[OMP_DONE:%.*]], label [[OMP_ARRAYMAP_BODY:%.*]]
// CHECK-60: omp.arraymap.body:
-// CHECK-60: [[OMP_ARRAYMAP_PTRCURRENT:%.*]] = phi ptr [ [[TMP2]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END25:%.*]] ]
+// CHECK-60: [[OMP_ARRAYMAP_PTRCURRENT:%.*]] = phi ptr [ [[TMP2]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END33:%.*]] ]
// CHECK-60: [[Z:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 1
// CHECK-60: [[S1P:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
+// CHECK-60: [[TMP15:%.*]] = load ptr, ptr [[S1P]], align 8
// CHECK-60: [[S1P1:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
-// CHECK-60: [[TMP15:%.*]] = load ptr, ptr [[S1P1]], align 8
-// CHECK-60: [[X:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP15]], i32 0, i32 0
+// CHECK-60: [[TMP16:%.*]] = load ptr, ptr [[S1P1]], align 8
+// CHECK-60: [[X:%.*]] = getelementptr inbounds nuw [[STRUCT_S1:%.*]], ptr [[TMP16]], i32 0, i32 0
// CHECK-60: [[S1P2:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
+// CHECK-60: [[TMP17:%.*]] = load ptr, ptr [[S1P2]], align 8
// CHECK-60: [[S1P3:%.*]] = getelementptr inbounds nuw [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
-// CHECK-60: [[TMP16:%.*]] = load ptr, ptr [[S1P3]], align 8
-// CHECK-60: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S1]], ptr [[TMP16]], i32 0, i32 1
-// CHECK-60: [[TMP17:%.*]] = getelementptr i32, ptr [[Z]], i32 1
-// CHECK-60: [[TMP18:%.*]] = ptrtoaddr ptr [[TMP17]] to i64
-// CHECK-60: [[TMP19:%.*]] = ptrtoaddr ptr [[S1P]] to i64
-// CHECK-60: [[TMP20:%.*]] = sub i64 [[TMP18]], [[TMP19]]
-// CHECK-60: [[TMP21:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP0]])
-// CHECK-60: [[TMP22:%.*]] = shl i64 [[TMP21]], 48
-// CHECK-60: [[TMP23:%.*]] = add nuw i64 0, [[TMP22]]
-// CHECK-60: [[TMP24:%.*]] = and i64 [[TMP4]], 3
-// CHECK-60: [[TMP25:%.*]] = icmp eq i64 [[TMP24]], 0
-// CHECK-60: br i1 [[TMP25]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]]
+// CHECK-60: [[TMP18:%.*]] = load ptr, ptr [[S1P3]], align 8
+// CHECK-60: [[Y:%.*]] = getelementptr inbounds nuw [[STRUCT_S1]], ptr [[TMP18]], i32 0, i32 1
+// CHECK-60: [[TMP19:%.*]] = getelementptr i32, ptr [[Y]], i32 1
+// CHECK-60: [[TMP20:%.*]] = ptrtoaddr ptr [[TMP19]] to i64
+// CHECK-60: [[TMP21:%.*]] = ptrtoaddr ptr [[X]] to i64
+// CHECK-60: [[TMP22:%.*]] = sub i64 [[TMP20]], [[TMP21]]
+// CHECK-60: [[TMP23:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP0]])
+// CHECK-60: [[TMP24:%.*]] = shl i64 [[TMP23]], 48
+// CHECK-60: [[TMP25:%.*]] = add nuw i64 3, [[TMP24]]
+// CHECK-60: [[TMP26:%.*]] = and i64 [[TMP4]], 3
+// CHECK-60: [[TMP27:%.*]] = icmp eq i64 [[TMP26]], 0
+// CHECK-60: br i1 [[TMP27]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]]
// CHECK-60: omp.type.alloc:
-// CHECK-60: [[TMP26:%.*]] = and i64 [[TMP23]], -4
+// CHECK-60: [[TMP28:%.*]] = and i64 [[TMP25]], -4
// CHECK-60: br label [[OMP_TYPE_END:%.*]]
// CHECK-60: omp.type.alloc.else:
-// CHECK-60: [[TMP27:%.*]] = icmp eq i64 [[TMP24]], 1
-// CHECK-60: br i1 [[TMP27]], label [[OMP_TYPE_TO:%.*]], label [[OMP_TYPE_TO_ELSE:%.*]]
+// CHECK-60: [[TMP29:%.*]] = icmp eq i64 [[TMP26]], 1
+// CHECK-60: br i1 [[TMP29]], label [[OMP_TYPE_TO:%.*]], label [[OMP_TYPE_TO_ELSE:%.*]]
// CHECK-60: omp.type.to:
-// CHECK-60: [[TMP28:%.*]] = and i64 [[TMP23]], -3
+// CHECK-60: [[TMP30:%.*]] = and i64 [[TMP25]], -3
// CHECK-60: br label [[OMP_TYPE_END]]
// CHECK-60: omp.type.to.else:
-// CHECK-60: [[TMP29:%.*]] = icmp eq i64 [[TMP24]], 2
-// CHECK-60: br i1 [[TMP29]], label [[OMP_TYPE_FROM:%.*]], label [[OMP_TYPE_END]]
+// CHECK-60: [[TMP31:%.*]] = icmp eq i64 [[TMP26]], 2
+// CHECK-60: br i1 [[TMP31]], label [[OMP_TYPE_FROM:%.*]], label [[OMP_TYPE_END]]
// CHECK-60: omp.type.from:
-// CHECK-60: [[TMP30:%.*]] = and i64 [[TMP23]], -2
+// CHECK-60: [[TMP32:%.*]] = and i64 [[TMP25]], -2
// CHECK-60: br label [[OMP_TYPE_END]]
// CHECK-60: omp.type.end:
-// CHECK-60: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP26]], [[OMP_TYPE_ALLOC]] ], [ [[TMP28]], [[OMP_TYPE_TO]] ], [ [[TMP30]], [[OMP_TYPE_FROM]] ], [ [[TMP23]], [[OMP_TYPE_TO_ELSE]] ]
-// CHECK-60: [[TMP31:%.*]] = and i64 [[TMP4]], 1036
-// CHECK-60: [[OMP_MAPTYPE_WITH_MODIFIERS:%.*]] = or i64 [[OMP_MAPTYPE]], [[TMP31]]
-// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[S1P]], i64 [[TMP20]], i64 [[OMP_MAPTYPE_WITH_MODIFIERS]], ptr null)
-// CHECK-60: [[TMP32:%.*]] = add nuw i64 281474976710659, [[TMP22]]
-// CHECK-60: [[TMP33:%.*]] = and i64 [[TMP4]], 3
-// CHECK-60: [[TMP34:%.*]] = icmp eq i64 [[TMP33]], 0
-// CHECK-60: br i1 [[TMP34]], label [[OMP_TYPE_ALLOC4:%.*]], label [[OMP_TYPE_ALLOC_ELSE5:%.*]]
+// CHECK-60: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP28]], [[OMP_TYPE_ALLOC]] ], [ [[TMP30]], [[OMP_TYPE_TO]] ], [ [[TMP32]], [[OMP_TYPE_FROM]] ], [ [[TMP25]], [[OMP_TYPE_TO_ELSE]] ]
+// CHECK-60: [[TMP33:%.*]] = and i64 [[TMP4]], 1036
+// CHECK-60: [[OMP_MAPTYPE_WITH_MODIFIERS:%.*]] = or i64 [[OMP_MAPTYPE]], [[TMP33]]
+// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[Z]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS]], ptr null)
+// CHECK-60: [[TMP34:%.*]] = and i64 [[TMP4]], 3
+// CHECK-60: [[TMP35:%.*]] = icmp eq i64 [[TMP34]], 0
+// CHECK-60: br i1 [[TMP35]], label [[OMP_TYPE_ALLOC4:%.*]], label [[OMP_TYPE_ALLOC_ELSE5:%.*]]
// CHECK-60: omp.type.alloc4:
-// CHECK-60: [[TMP35:%.*]] = and i64 [[TMP32]], -4
// CHECK-60: br label [[OMP_TYPE_END9:%.*]]
// CHECK-60: omp.type.alloc.else5:
-// CHECK-60: [[TMP36:%.*]] = icmp eq i64 [[TMP33]], 1
+// CHECK-60: [[TMP36:%.*]] = icmp eq i64 [[TMP34]], 1
// CHECK-60: br i1 [[TMP36]], label [[OMP_TYPE_TO6:%.*]], label [[OMP_TYPE_TO_ELSE7:%.*]]
// CHECK-60: omp.type.to6:
-// CHECK-60: [[TMP37:%.*]] = and i64 [[TMP32]], -3
// CHECK-60: br label [[OMP_TYPE_END9]]
// CHECK-60: omp.type.to.else7:
-// CHECK-60: [[TMP38:%.*]] = icmp eq i64 [[TMP33]], 2
-// CHECK-60: br i1 [[TMP38]], label [[OMP_TYPE_FROM8:%.*]], label [[OMP_TYPE_END9]]
+// CHECK-60: [[TMP37:%.*]] = icmp eq i64 [[TMP34]], 2
+// CHECK-60: br i1 [[TMP37]], label [[OMP_TYPE_FROM8:%.*]], label [[OMP_TYPE_END9]]
// CHECK-60: omp.type.from8:
-// CHECK-60: [[TMP39:%.*]] = and i64 [[TMP32]], -2
// CHECK-60: br label [[OMP_TYPE_END9]]
// CHECK-60: omp.type.end9:
-// CHECK-60: [[OMP_MAPTYPE10:%.*]] = phi i64 [ [[TMP35]], [[OMP_TYPE_ALLOC4]] ], [ [[TMP37]], [[OMP_TYPE_TO6]] ], [ [[TMP39]], [[OMP_TYPE_FROM8]] ], [ [[TMP32]], [[OMP_TYPE_TO_ELSE7]] ]
-// CHECK-60: [[TMP40:%.*]] = and i64 [[TMP4]], 1036
-// CHECK-60: [[OMP_MAPTYPE_WITH_MODIFIERS11:%.*]] = or i64 [[OMP_MAPTYPE10]], [[TMP40]]
-// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[Z]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS11]], ptr null)
-// CHECK-60: [[TMP41:%.*]] = add nuw i64 281474976710675, [[TMP22]]
-// CHECK-60: [[TMP42:%.*]] = and i64 [[TMP4]], 3
-// CHECK-60: [[TMP43:%.*]] = icmp eq i64 [[TMP42]], 0
-// CHECK-60: br i1 [[TMP43]], label [[OMP_TYPE_ALLOC12:%.*]], label [[OMP_TYPE_ALLOC_ELSE13:%.*]]
+// CHECK-60: [[OMP_MAPTYPE10:%.*]] = phi i64 [ 0, [[OMP_TYPE_ALLOC4]] ], [ 0, [[OMP_TYPE_TO6]] ], [ 0, [[OMP_TYPE_FROM8]] ], [ 0, [[OMP_TYPE_TO_ELSE7]] ]
+// CHECK-60: [[TMP38:%.*]] = and i64 [[TMP4]], 1036
+// CHECK-60: [[OMP_MAPTYPE_WITH_MODIFIERS11:%.*]] = or i64 [[OMP_MAPTYPE10]], [[TMP38]]
+// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP15]], ptr [[X]], i64 [[TMP22]], i64 [[OMP_MAPTYPE_WITH_MODIFIERS11]], ptr null)
+// CHECK-60: [[TMP39:%.*]] = add nuw i64 562949953421315, [[TMP24]]
+// CHECK-60: [[TMP40:%.*]] = and i64 [[TMP4]], 3
+// CHECK-60: [[TMP41:%.*]] = icmp eq i64 [[TMP40]], 0
+// CHECK-60: br i1 [[TMP41]], label [[OMP_TYPE_ALLOC12:%.*]], label [[OMP_TYPE_ALLOC_ELSE13:%.*]]
// CHECK-60: omp.type.alloc12:
-// CHECK-60: [[TMP44:%.*]] = and i64 [[TMP41]], -4
+// CHECK-60: [[TMP42:%.*]] = and i64 [[TMP39]], -4
// CHECK-60: br label [[OMP_TYPE_END17:%.*]]
// CHECK-60: omp.type.alloc.else13:
-// CHECK-60: [[TMP45:%.*]] = icmp eq i64 [[TMP42]], 1
-// CHECK-60: br i1 [[TMP45]], label [[OMP_TYPE_TO14:%.*]], label [[OMP_TYPE_TO_ELSE15:%.*]]
+// CHECK-60: [[TMP43:%.*]] = icmp eq i64 [[TMP40]], 1
+// CHECK-60: br i1 [[TMP43]], label [[OMP_TYPE_TO14:%.*]], label [[OMP_TYPE_TO_ELSE15:%.*]]
// CHECK-60: omp.type.to14:
-// CHECK-60: [[TMP46:%.*]] = and i64 [[TMP41]], -3
+// CHECK-60: [[TMP44:%.*]] = and i64 [[TMP39]], -3
// CHECK-60: br label [[OMP_TYPE_END17]]
// CHECK-60: omp.type.to.else15:
-// CHECK-60: [[TMP47:%.*]] = icmp eq i64 [[TMP42]], 2
-// CHECK-60: br i1 [[TMP47]], label [[OMP_TYPE_FROM16:%.*]], label [[OMP_TYPE_END17]]
+// CHECK-60: [[TMP45:%.*]] = icmp eq i64 [[TMP40]], 2
+// CHECK-60: br i1 [[TMP45]], label [[OMP_TYPE_FROM16:%.*]], label [[OMP_TYPE_END17]]
// CHECK-60: omp.type.from16:
-// CHECK-60: [[TMP48:%.*]] = and i64 [[TMP41]], -2
+// CHECK-60: [[TMP46:%.*]] = and i64 [[TMP39]], -2
// CHECK-60: br label [[OMP_TYPE_END17]]
// CHECK-60: omp.type.end17:
-// CHECK-60: [[OMP_MAPTYPE18:%.*]] = phi i64 [ [[TMP44]], [[OMP_TYPE_ALLOC12]] ], [ [[TMP46]], [[OMP_TYPE_TO14]] ], [ [[TMP48]], [[OMP_TYPE_FROM16]] ], [ [[TMP41]], [[OMP_TYPE_TO_ELSE15]] ]
-// CHECK-60: [[TMP49:%.*]] = and i64 [[TMP4]], 1036
-// CHECK-60: [[OMP_MAPTYPE_WITH_MODIFIERS19:%.*]] = or i64 [[OMP_MAPTYPE18]], [[TMP49]]
-// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[S1P]], ptr [[X]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS19]], ptr null)
-// CHECK-60: [[TMP50:%.*]] = add nuw i64 281474976710675, [[TMP22]]
-// CHECK-60: [[TMP51:%.*]] = and i64 [[TMP4]], 3
-// CHECK-60: [[TMP52:%.*]] = icmp eq i64 [[TMP51]], 0
-// CHECK-60: br i1 [[TMP52]], label [[OMP_TYPE_ALLOC20:%.*]], label [[OMP_TYPE_ALLOC_ELSE21:%.*]]
+// CHECK-60: [[OMP_MAPTYPE18:%.*]] = phi i64 [ [[TMP42]], [[OMP_TYPE_ALLOC12]] ], [ [[TMP44]], [[OMP_TYPE_TO14]] ], [ [[TMP46]], [[OMP_TYPE_FROM16]] ], [ [[TMP39]], [[OMP_TYPE_TO_ELSE15]] ]
+// CHECK-60: [[TMP47:%.*]] = and i64 [[TMP4]], 1036
+// CHECK-60: [[OMP_MAPTYPE_WITH_MODIFIERS19:%.*]] = or i64 [[OMP_MAPTYPE18]], [[TMP47]]
+// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP15]], ptr [[X]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS19]], ptr null)
+// CHECK-60: [[TMP48:%.*]] = add nuw i64 562949953421315, [[TMP24]]
+// CHECK-60: [[TMP49:%.*]] = and i64 [[TMP4]], 3
+// CHECK-60: [[TMP50:%.*]] = icmp eq i64 [[TMP49]], 0
+// CHECK-60: br i1 [[TMP50]], label [[OMP_TYPE_ALLOC20:%.*]], label [[OMP_TYPE_ALLOC_ELSE21:%.*]]
// CHECK-60: omp.type.alloc20:
-// CHECK-60: [[TMP53:%.*]] = and i64 [[TMP50]], -4
-// CHECK-60: br label [[OMP_TYPE_END25]]
+// CHECK-60: [[TMP51:%.*]] = and i64 [[TMP48]], -4
+// CHECK-60: br label [[OMP_TYPE_END25:%.*]]
// CHECK-60: omp.type.alloc.else21:
-// CHECK-60: [[TMP54:%.*]] = icmp eq i64 [[TMP51]], 1
-// CHECK-60: br i1 [[TMP54]], label [[OMP_TYPE_TO22:%.*]], label [[OMP_TYPE_TO_ELSE23:%.*]]
+// CHECK-60: [[TMP52:%.*]] = icmp eq i64 [[TMP49]], 1
+// CHECK-60: br i1 [[TMP52]], label [[OMP_TYPE_TO22:%.*]], label [[OMP_TYPE_TO_ELSE23:%.*]]
// CHECK-60: omp.type.to22:
-// CHECK-60: [[TMP55:%.*]] = and i64 [[TMP50]], -3
+// CHECK-60: [[TMP53:%.*]] = and i64 [[TMP48]], -3
// CHECK-60: br label [[OMP_TYPE_END25]]
// CHECK-60: omp.type.to.else23:
-// CHECK-60: [[TMP56:%.*]] = icmp eq i64 [[TMP51]], 2
-// CHECK-60: br i1 [[TMP56]], label [[OMP_TYPE_FROM24:%.*]], label [[OMP_TYPE_END25]]
+// CHECK-60: [[TMP54:%.*]] = icmp eq i64 [[TMP49]], 2
+// CHECK-60: br i1 [[TMP54]], label [[OMP_TYPE_FROM24:%.*]], label [[OMP_TYPE_END25]]
// CHECK-60: omp.type.from24:
-// CHECK-60: [[TMP57:%.*]] = and i64 [[TMP50]], -2
+// CHECK-60: [[TMP55:%.*]] = and i64 [[TMP48]], -2
// CHECK-60: br label [[OMP_TYPE_END25]]
// CHECK-60: omp.type.end25:
-// CHECK-60: [[OMP_MAPTYPE26:%.*]] = phi i64 [ [[TMP53]], [[OMP_TYPE_ALLOC20]] ], [ [[TMP55]], [[OMP_TYPE_TO22]] ], [ [[TMP57]], [[OMP_TYPE_FROM24]] ], [ [[TMP50]], [[OMP_TYPE_TO_ELSE23]] ]
-// CHECK-60: [[TMP58:%.*]] = and i64 [[TMP4]], 1036
-// CHECK-60: [[OMP_MAPTYPE_WITH_MODIFIERS27:%.*]] = or i64 [[OMP_MAPTYPE26]], [[TMP58]]
-// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[S1P2]], ptr [[Y]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS27]], ptr null)
+// CHECK-60: [[OMP_MAPTYPE26:%.*]] = phi i64 [ [[TMP51]], [[OMP_TYPE_ALLOC20]] ], [ [[TMP53]], [[OMP_TYPE_TO22]] ], [ [[TMP55]], [[OMP_TYPE_FROM24]] ], [ [[TMP48]], [[OMP_TYPE_TO_ELSE23]] ]
+// CHECK-60: [[TMP56:%.*]] = and i64 [[TMP4]], 1036
+// CHECK-60: [[OMP_MAPTYPE_WITH_MODIFIERS27:%.*]] = or i64 [[OMP_MAPTYPE26]], [[TMP56]]
+// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP17]], ptr [[Y]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS27]], ptr null)
+// CHECK-60: [[TMP57:%.*]] = and i64 [[TMP4]], 3
+// CHECK-60: [[TMP58:%.*]] = icmp eq i64 [[TMP57]], 0
+// CHECK-60: br i1 [[TMP58]], label [[OMP_TYPE_ALLOC28:%.*]], label [[OMP_TYPE_ALLOC_ELSE29:%.*]]
+// CHECK-60: omp.type.alloc28:
+// CHECK-60: br label [[OMP_TYPE_END33]]
+// CHECK-60: omp.type.alloc.else29:
+// CHECK-60: [[TMP59:%.*]] = icmp eq i64 [[TMP57]], 1
+// CHECK-60: br i1 [[TMP59]], label [[OMP_TYPE_TO30:%.*]], label [[OMP_TYPE_TO_ELSE31:%.*]]
+// CHECK-60: omp.type.to30:
+// CHECK-60: br label [[OMP_TYPE_END33]]
+// CHECK-60: omp.type.to.else31:
+// CHECK-60: [[TMP60:%.*]] = icmp eq i64 [[TMP57]], 2
+// CHECK-60: br i1 [[TMP60]], label [[OMP_TYPE_FROM32:%.*]], label [[OMP_TYPE_END33]]
+// CHECK-60: omp.type.from32:
+// CHECK-60: br label [[OMP_TYPE_END33]]
+// CHECK-60: omp.type.end33:
+// CHECK-60: [[OMP_MAPTYPE34:%.*]] = phi i64 [ 16384, [[OMP_TYPE_ALLOC28]] ], [ 16384, [[OMP_TYPE_TO30]] ], [ 16384, [[OMP_TYPE_FROM32]] ], [ 16384, [[OMP_TYPE_TO_ELSE31]] ]
+// CHECK-60: [[TMP61:%.*]] = and i64 [[TMP4]], 1036
+// CHECK-60: [[OMP_MAPTYPE_WITH_MODIFIERS35:%.*]] = or i64 [[OMP_MAPTYPE34]], [[TMP61]]
+// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[S1P2]], ptr [[X]], i64 8, i64 [[OMP_MAPTYPE34]], ptr null)
// CHECK-60: [[OMP_ARRAYMAP_NEXT]] = getelementptr [[STRUCT_S2]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 1
// CHECK-60: [[OMP_ARRAYMAP_ISDONE:%.*]] = icmp eq ptr [[OMP_ARRAYMAP_NEXT]], [[TMP7]]
// CHECK-60: br i1 [[OMP_ARRAYMAP_ISDONE]], label [[OMP_ARRAYMAP_EXIT:%.*]], label [[OMP_ARRAYMAP_BODY]]
// CHECK-60: omp.arraymap.exit:
-// CHECK-60: [[OMP_ARRAYINIT_ISARRAY28:%.*]] = icmp sgt i64 [[TMP6]], 1
-// CHECK-60: [[TMP59:%.*]] = and i64 [[TMP4]], 8
-// CHECK-60: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP59]], 0
-// CHECK-60: [[TMP60:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY28]], [[DOTOMP_ARRAY__DEL__DELETE]]
-// CHECK-60: br i1 [[TMP60]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]]
+// CHECK-60: [[OMP_ARRAYINIT_ISARRAY36:%.*]] = icmp sgt i64 [[TMP6]], 1
+// CHECK-60: [[TMP62:%.*]] = and i64 [[TMP4]], 8
+// CHECK-60: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP62]], 0
+// CHECK-60: [[TMP63:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY36]], [[DOTOMP_ARRAY__DEL__DELETE]]
+// CHECK-60: br i1 [[TMP63]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]]
// CHECK-60: .omp.array..del:
-// CHECK-60: [[TMP61:%.*]] = mul nuw i64 [[TMP6]], 16
-// CHECK-60: [[TMP62:%.*]] = and i64 [[TMP4]], -4
-// CHECK-60: [[TMP63:%.*]] = or i64 [[TMP62]], 512
-// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP61]], i64 [[TMP63]], ptr [[TMP5]])
+// CHECK-60: [[TMP64:%.*]] = mul nuw i64 [[TMP6]], 16
+// CHECK-60: [[TMP65:%.*]] = and i64 [[TMP4]], -4
+// CHECK-60: [[TMP66:%.*]] = or i64 [[TMP65]], 512
+// CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP64]], i64 [[TMP66]], ptr [[TMP5]])
// CHECK-60: br label [[OMP_DONE]]
// CHECK-60: omp.done:
// CHECK-60: ret void
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
index d4798f5a92dc6..2fd98feea0c33 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
@@ -2923,6 +2923,7 @@ class OpenMPIRBuilder {
using MapNamesArrayTy = SmallVector<Constant *, 4>;
using MapDimArrayTy = SmallVector<uint64_t, 4>;
using MapNonContiguousArrayTy = SmallVector<MapValuesArrayTy, 4>;
+ using MapHasAttachPtrArrayTy = SmallVector<bool, 4>;
/// This structure contains combined information generated for mappable
/// clauses, including base pointers, pointers, sizes, map types, user-defined
@@ -2941,6 +2942,9 @@ class OpenMPIRBuilder {
MapValuesArrayTy Sizes;
MapFlagsArrayTy Types;
MapNamesArrayTy Names;
+ /// True for entries that have an attach ptr, and thus an accompanying
+ /// ATTACH entry linking that ptr to its ptee.
+ MapHasAttachPtrArrayTy HasAttachPtr;
StructNonContiguousInfo NonContigInfo;
/// Append arrays in \a CurInfo.
@@ -2953,6 +2957,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());
+ HasAttachPtr.append(CurInfo.HasAttachPtr.begin(),
+ CurInfo.HasAttachPtr.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 9704f68df4ec7..3df035f75746c 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -10519,19 +10519,66 @@ 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]));
+ auto RawType =
+ static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>(
+ Info->Types[I]);
+ constexpr uint64_t MemberOfMask =
+ static_cast<uint64_t>(OpenMPOffloadMappingFlags::OMP_MAP_MEMBER_OF);
+ constexpr uint64_t AttachBit =
+ static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>(
+ OpenMPOffloadMappingFlags::OMP_MAP_ATTACH);
+
+ // 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])
+ // 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 (**)
+ //
+ // 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, 4, MEMBER_OF(N+2)|TO|FROM (***)
+ // &arr[i].s1p[0], &arr[i].s1p->y, 4, 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 of the following kinds do NOT receive a new outer MEMBER_OF
+ // linking them to the parent struct:
+ //
+ // * (*) Entries with HasAttachPtr: they represent pointee data that
+ // occupies a different storage block than the struct being mapped, so
+ // they are not a member of it.
+ // * (**) ATTACH entries: they are not a member of anything — they just
+ // link a ptr to its ptee.
+ // * All entries when PreserveMemberOfFlags is set (the Flang/MLIR path):
+ // its pre-shaped entries already carry their final MEMBER_OF bits.
+ // TODO: set HasAttachPtr from Flang for entries whose storage is the
+ // pointee's (e.g. s%p(0:10)) and drop PreserveMemberOfFlags in favor of
+ // it.
+ //
+ // (***) If such an entry already has its own MEMBER_OF bits (e.g. the
+ // s1p->x/y entries above), those bits are still shifted by N.
Value *MemberMapType;
- if (PreserveMemberOfFlags) {
- constexpr uint64_t MemberOfMask =
- static_cast<uint64_t>(OpenMPOffloadMappingFlags::OMP_MAP_MEMBER_OF);
- uint64_t OrigFlags =
- static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>(
- Info->Types[I]);
- bool HasMemberOf = (OrigFlags & MemberOfMask) != 0;
- if (HasMemberOf)
+ if (PreserveMemberOfFlags || (RawType & AttachBit) ||
+ Info->HasAttachPtr[I]) {
+ if (RawType & MemberOfMask)
MemberMapType = Builder.CreateNUWAdd(OriMapType, ShiftedPreviousSize);
else
MemberMapType = OriMapType;
@@ -10642,12 +10689,6 @@ 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) have no meaning for an ATTACH entry.
- auto RawType =
- static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>(
- Info->Types[I]);
- constexpr uint64_t AttachBit =
- static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>(
- OpenMPOffloadMappingFlags::OMP_MAP_ATTACH);
Value *FinalMapType =
(RawType & AttachBit) ? CurMapType : CurMapTypeWithModifiers;
diff --git a/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp b/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp
index b96103c82c185..14b690fe34ec2 100644
--- a/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp
+++ b/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp
@@ -8312,6 +8312,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.HasAttachPtr.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 1c50ff192c3d5..86d1cf4f740ed 100644
--- a/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp
+++ b/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp
@@ -6697,6 +6697,8 @@ static void collectMapDataFromMapOperands(
builder, moduleTranslation));
mapData.MapClause.push_back(mapOp.getOperation());
mapData.Types.push_back(convertClauseMapFlags(mapOp.getMapType()));
+ // TODO: set HasAttachPtr from Flang for pointee-storage entries.
+ mapData.HasAttachPtr.push_back(false);
mapData.Names.push_back(LLVM::createMappingInformation(
mapOp.getLoc(), *moduleTranslation.getOpenMPBuilder()));
mapData.DevicePointers.push_back(llvm::OpenMPIRBuilder::DeviceInfoTy::None);
@@ -6765,6 +6767,8 @@ static void collectMapDataFromMapOperands(
mapData.MapClause.push_back(mapOp.getOperation());
mapData.Types.push_back(
llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_RETURN_PARAM);
+ // TODO: set HasAttachPtr from Flang for pointee-storage entries.
+ mapData.HasAttachPtr.push_back(false);
mapData.Names.push_back(LLVM::createMappingInformation(
mapOp.getLoc(), *moduleTranslation.getOpenMPBuilder()));
mapData.DevicePointers.push_back(devInfoTy);
@@ -6805,6 +6809,8 @@ 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);
+ // TODO: set HasAttachPtr from Flang for pointee-storage entries.
+ mapData.HasAttachPtr.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.
@@ -6821,6 +6827,8 @@ static void collectMapDataFromMapOperands(
mapData.Types.push_back(
isDevicePtr ? mapType
: llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_LITERAL);
+ // TODO: set HasAttachPtr from Flang for pointee-storage entries.
+ mapData.HasAttachPtr.push_back(false);
mapData.Mappers.push_back(nullptr);
}
mapData.Names.push_back(LLVM::createMappingInformation(
@@ -7126,6 +7134,8 @@ processIndividualMap(llvm::IRBuilderBase &builder,
combinedInfo.Mappers.emplace_back(mapData.Mappers[mapDataIdx]);
combinedInfo.Names.emplace_back(mapData.Names[mapDataIdx]);
combinedInfo.Types.emplace_back(mapFlag);
+ // TODO: set HasAttachPtr from Flang for pointee-storage entries.
+ combinedInfo.HasAttachPtr.emplace_back(false);
combinedInfo.Sizes.emplace_back(
isPtrTy ? builder.CreateSelect(
builder.CreateIsNull(mapData.Pointers[mapDataIdx]),
@@ -7189,6 +7199,8 @@ static void mapParentWithMembers(
}
combinedInfo.Types.emplace_back(baseFlag);
+ // TODO: set HasAttachPtr from Flang for pointee-storage entries.
+ combinedInfo.HasAttachPtr.emplace_back(false);
combinedInfo.DevicePointers.emplace_back(
mapData.DevicePointers[mapDataIndex]);
// Only attach the mapper to the base entry when we are mapping the whole
@@ -7289,6 +7301,8 @@ static void mapParentWithMembers(
if (targetDirective == TargetDirectiveEnumTy::TargetUpdate || hasMapClose ||
overlapIdxs.size() == 1) {
combinedInfo.Types.emplace_back(mapFlag);
+ // TODO: set HasAttachPtr from Flang for pointee-storage entries.
+ combinedInfo.HasAttachPtr.emplace_back(false);
combinedInfo.DevicePointers.emplace_back(
mapData.DevicePointers[mapDataIndex]);
combinedInfo.Names.emplace_back(LLVM::createMappingInformation(
@@ -7329,6 +7343,8 @@ static void mapParentWithMembers(
auto isPtrMap = checkIfPointerMap(
llvm::cast<omp::MapInfoOp>(mapData.MapClause[mapDataOverlapIdx]));
combinedInfo.Types.emplace_back(mapFlag);
+ // TODO: set HasAttachPtr from Flang for pointee-storage entries.
+ combinedInfo.HasAttachPtr.emplace_back(false);
combinedInfo.DevicePointers.emplace_back(
llvm::OpenMPIRBuilder::DeviceInfoTy::None);
combinedInfo.Names.emplace_back(LLVM::createMappingInformation(
@@ -7357,6 +7373,8 @@ static void mapParentWithMembers(
}
combinedInfo.Types.emplace_back(mapFlag);
+ // TODO: set HasAttachPtr from Flang for pointee-storage entries.
+ combinedInfo.HasAttachPtr.emplace_back(false);
combinedInfo.DevicePointers.emplace_back(
llvm::OpenMPIRBuilder::DeviceInfoTy::None);
combinedInfo.Names.emplace_back(LLVM::createMappingInformation(
@@ -8998,6 +9016,8 @@ convertOmpTarget(Operation &opInst, llvm::IRBuilderBase &builder,
combinedInfos.Types.push_back(
llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_TARGET_PARAM |
llvm::omp::OpenMPOffloadMappingFlags::OMP_MAP_LITERAL);
+ // TODO: set HasAttachPtr from Flang for pointee-storage entries.
+ combinedInfos.HasAttachPtr.push_back(false);
if (!combinedInfos.Names.empty())
combinedInfos.Names.push_back(nullPtr);
combinedInfos.Mappers.push_back(nullptr);
diff --git a/offload/test/mapping/mapper_enter_data_always_present_ptee.c b/offload/test/mapping/mapper_enter_data_always_present_ptee.c
index 9fd3a6bf0d794..810ec89eb8eb9 100644
--- a/offload/test/mapping/mapper_enter_data_always_present_ptee.c
+++ b/offload/test/mapping/mapper_enter_data_always_present_ptee.c
@@ -2,31 +2,27 @@
// "target enter data map(always, present, to : s)" clause instead of a
// "target update to(present : s)" motion clause. Both invoke the mapper; this
// checks present propagation to the pointee is consistent across the two paths.
-//
-// FIXME: this test currently run-fails at every version/bounds combination.
-// The mapper maps the struct member (s.y) with a combined entry whose size does
-// not match the member's own storage, so the map clause aborts with an
-// "explicit extension not allowed" error before the present modifier is ever
-// considered. This is fixed once the mapper emits attach-style maps for pointer
-// members (so the member and pointee occupy separate, correctly-sized entries).
-//
-// EXPECTED final state:
-// inbounds, 5.2 and 6.0: run succeeds, prints "333 333".
-// out-of-bounds (s.p[0:20] over the mapped x[0:10]):
-// 5.2: succeeds (present is not applied to the pointee before 6.0).
-// 6.0: run-fails; the failure should be the 'present' map-type-modifier
-// check on s.p[0:20] (an accompanying "explicit extension" message is
-// incidental -- for a map clause it is user error to map 20 elements
-// when only 10 are present).
+// Inbounds: the pointee region is fully present; the run succeeds at every
+// OpenMP version.
// RUN: %libomptarget-compile-generic -fopenmp-version=52
-// RUN: %libomptarget-run-fail-generic 2>&1 | %fcheck-generic
+// RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic
// RUN: %libomptarget-compile-generic -fopenmp-version=60
-// RUN: %libomptarget-run-fail-generic 2>&1 | %fcheck-generic
+// RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic
+
+// Out-of-bounds: s.p[0:20] extends beyond the mapped region x[0:10]. Because a
+// map clause performs an actual mapping, requesting 20 elements when only 10
+// are present is an error, so the run fails with an "explicit extension"
+// diagnostic at every version.
+// FIXME: at OpenMP 6.0 the present modifier is also propagated to the pointee,
+// so the failure should additionally report the 'present' map-type-modifier
+// check on s.p[0:20]. The extension diagnostic currently fires first.
// RUN: %libomptarget-compile-generic -fopenmp-version=52 -DOUT_OF_BOUNDS
-// RUN: %libomptarget-run-fail-generic 2>&1 | %fcheck-generic
+// RUN: %libomptarget-run-fail-generic 2>&1 \
+// RUN: | %fcheck-generic --check-prefix=CHECK-OOB
// RUN: %libomptarget-compile-generic -fopenmp-version=60 -DOUT_OF_BOUNDS
-// RUN: %libomptarget-run-fail-generic 2>&1 | %fcheck-generic
+// RUN: %libomptarget-run-fail-generic 2>&1 \
+// RUN: | %fcheck-generic --check-prefix=CHECK-OOB
#include <stdio.h>
@@ -63,14 +59,12 @@ int main() {
fprintf(stderr, "addr=%p, size=%zu\n", &s.p[0], 20 * sizeof(s.p[0]));
- // FIXME: the map clause aborts here with an "explicit extension not allowed"
- // error on the mapper's combined member entry, at every version. Fixed once
- // the mapper uses attach-style maps for pointer members.
- // CHECK: explicit extension not allowed
+ // CHECK-OOB: explicit extension not allowed
#pragma omp target data map(from : s.y, x)
{
f1();
}
+ // CHECK: 333 333
printf("%d %d\n", x[0], s.y);
}
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 6dea1acfdd9e3..fb3abc68159ce 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
@@ -1,15 +1,7 @@
// Check that it's ok to first map a member of a struct and its pointee, and
// then do a map(present) on a mapper that maps them internally.
-//
-// FIXME: This currently run-fails, because the mapper does not yet emit
-// attach-style maps for the pointee: the combined entry over the whole struct
-// s1 (40016 bytes) triggers an "explicit extension not allowed" error against
-// the 4-byte device allocation of s1.x. Once attach-style maps are emitted for
-// the pointee, the present check should pass and the run should complete the
-// present/delete sequence below.
-// RUN: %libomptarget-compile-generic
-// RUN: %libomptarget-run-fail-generic 2>&1 \
-// RUN: | %fcheck-generic --check-prefix=CHECK
+
+// RUN: %libomptarget-compile-run-and-check-generic
#include <omp.h>
#include <stdio.h>
@@ -35,24 +27,20 @@ int main() {
s1.p = (int *)&x;
#pragma omp target enter data map(alloc : s1.x, s1.p[0 : 10])
- // EXPECTED: After mapping
- print_status(&s1.x, "x"); // EXPECTED: x is present
- print_status(&s1.dummy, "dummy"); // EXPECTED: dummy is not present
- print_status(&s1.p, "p"); // EXPECTED: p is not present
- print_status(&s1.p[0], "p[0]"); // EXPECTED: p[0] is present
-
- // This present check currently fails (explicit extension); once attach-style
- // maps are emitted for the pointee, it should pass.
- // clang-format off
- // CHECK: omptarget message: explicit extension not allowed
- // CHECK: omptarget fatal error 1: failure of target construct while offloading is mandatory
- // clang-format on
+ 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");
+
+ // This present check should pass.
#pragma omp target enter data map(present, alloc : s1)
#pragma omp target exit data map(delete : s1)
- // EXPECTED: After deleting
- print_status(&s1.x, "x"); // EXPECTED: x is not present
- print_status(&s1.dummy, "dummy"); // EXPECTED: dummy is not present
- print_status(&s1.p, "p"); // EXPECTED: p is not present
- print_status(&s1.p[0], "p[0]"); // EXPECTED: p[0] is not present
+ 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
index 422bfa484c2dc..bb8b6aa3c3e76 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
@@ -1,22 +1,19 @@
// The mapper maps a struct member (s.x) and a pointee (s.p[0:10]). We pre-map
// only s.x, then do map(present) on the mapper. The pointee s.p[0:10] is not
// present, so once PRESENT is propagated to the pointee (a follow-on, at OpenMP
-// >= 6.0) the check must fail; at <= 5.2 present is not propagated, so it would
-// pass.
+// >= 6.0) the check must fail; at <= 5.2 present is not propagated, so it
+// passes.
//
-// FIXME: This currently run-fails at BOTH versions, because the mapper does not
-// yet emit attach-style maps for the pointee: the combined entry over the whole
-// struct s1 (40016 bytes) triggers an "explicit extension not allowed" error
-// against the 4-byte device allocation of s1.x. Once attach-style maps are
-// emitted for the pointee:
+// FIXME: PRESENT is not propagated to the pointee yet, so the run currently
+// completes ("done") at BOTH versions. Once it is propagated:
// EXPECTED (5.2): the run completes ("done").
// EXPECTED (6.0): the present check fails for the absent pointee s1.p[0:10].
// RUN: %libomptarget-compile-generic -fopenmp-version=52
-// RUN: %libomptarget-run-fail-generic 2>&1 \
-// RUN: | %fcheck-generic --check-prefixes=CHECK
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic --check-prefixes=CHECK,CHECK-52
// RUN: %libomptarget-compile-generic -fopenmp-version=60
-// RUN: %libomptarget-run-fail-generic 2>&1 \
-// RUN: | %fcheck-generic --check-prefixes=CHECK
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic --check-prefixes=CHECK,CHECK-60
#include <omp.h>
#include <stdio.h>
@@ -41,22 +38,21 @@ void print_status(void *p, const char *name) {
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)
- print_status(&s1.x, "x"); // EXPECTED: x is present
- print_status(&s1.dummy, "dummy"); // EXPECTED: dummy is not present
- print_status(&s1.p, "p"); // EXPECTED: p is not present
- print_status(&s1.p[0], "p[0]"); // EXPECTED: p[0] is not 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 not present
#pragma omp target enter data map(present, alloc : s1)
- // Once attach-style maps are emitted for the pointee, at 5.2 the run
- // completes past this point; at 6.0 the present check on the absent pointee
- // s1.p[0:10] fails here.
- // clang-format off
- // CHECK: omptarget message: explicit extension not allowed
- // CHECK: omptarget fatal error 1: failure of target construct while offloading is mandatory
- // clang-format on
+ // Once PRESENT is propagated to the pointee, at 5.2 the run completes past
+ // this point; at 6.0 the present check on the absent pointee s1.p[0:10]
+ // fails here.
+ // CHECK-52: done
+ // CHECK-60: done
fprintf(stderr, "done\n");
}
diff --git a/offload/test/mapping/mapper_map_present_ptee.c b/offload/test/mapping/mapper_map_present_ptee.c
index 065bbcb62e28f..e30663c8c2943 100644
--- a/offload/test/mapping/mapper_map_present_ptee.c
+++ b/offload/test/mapping/mapper_map_present_ptee.c
@@ -1,21 +1,12 @@
// A user-defined mapper maps s.y and, with present written directly in the
// mapper's own clause, the pointee s.p. present in the mapper clause applies at
// every OpenMP version -- this is not the outer-clause propagation that a
-// follow-on gates on the version.
+// follow-on gates on the version, so both 5.2 and 6.0 behave the same here.
//
-// FIXME: This currently run-fails at BOTH the inbounds and out-of-bounds cases,
-// because the mapper does not yet emit attach-style maps for the pointee: the
-// present check runs against the whole struct s (16 bytes), which is not
-// present. Once attach-style maps are emitted for the pointee:
-// EXPECTED (inbounds): the present check passes and the run prints "333 333".
-// EXPECTED (out-of-bounds): the present check fails only for the pointee
-// region s.p[0:20] that is not present.
-// RUN: %libomptarget-compile-generic
-// RUN: %libomptarget-run-fail-generic 2>&1 \
-// RUN: | %fcheck-generic --check-prefix=CHECK
+// RUN: %libomptarget-compile-run-and-check-generic
// RUN: %libomptarget-compile-generic -DOUT_OF_BOUNDS
// RUN: %libomptarget-run-fail-generic 2>&1 \
-// RUN: | %fcheck-generic --check-prefix=CHECK
+// RUN: | %fcheck-generic --check-prefix=CHECK-OOB
#include <stdio.h>
@@ -26,6 +17,7 @@ typedef struct {
} S;
#ifdef OUT_OF_BOUNDS
+// s.p[0:20] extends beyond the mapped region x[0:10]; the present check fails.
#pragma omp declare mapper(S s) map(s.y) map(present, tofrom : s.p[0 : 20])
#else
#pragma omp declare mapper(S s) map(s.y) map(present, tofrom : s.p[0 : 2])
@@ -33,7 +25,7 @@ typedef struct {
S s;
void f1() {
- // The mapper runs here; the present check currently fails here at both cases.
+ // The mapper runs here; for OUT_OF_BOUNDS the present check fails here.
#pragma omp target update to(s)
#pragma omp target data use_device_addr(s, x)
@@ -49,6 +41,7 @@ int main() {
s.y = 111;
s.p = &x[0];
+ // CHECK-OOB: addr=0x[[#%x,HOST_ADDR:]], size=[[#%u,SIZE:]]
fprintf(stderr, "addr=%p, size=%zu\n", &s.p[0], 20 * sizeof(s.p[0]));
#pragma omp target data map(from : s.y, x)
@@ -56,10 +49,9 @@ int main() {
f1();
}
- // EXPECTED (inbounds): 333 333
- fprintf(stderr, "%d %d\n", x[0], s.y);
+ printf("%d %d\n", x[0], s.y); // CHECK: 333 333
// clang-format off
- // CHECK: omptarget message: device mapping required by 'present' motion modifier does not exist for host address
- // CHECK: omptarget fatal error 1: failure of target construct while offloading is mandatory
+ // CHECK-OOB: omptarget message: device mapping required by 'present' motion modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes)
+ // CHECK-OOB: omptarget fatal error 1: failure of target construct while offloading is mandatory
// clang-format on
}
diff --git a/offload/test/mapping/mapper_map_ptee_only.c b/offload/test/mapping/mapper_map_ptee_only.c
index 6e9904bb1917e..3944545b6e23c 100644
--- a/offload/test/mapping/mapper_map_ptee_only.c
+++ b/offload/test/mapping/mapper_map_ptee_only.c
@@ -28,16 +28,10 @@ int main() {
#pragma omp target enter data map(alloc : s1)
printf("After mapping\n");
- print_status(&s1.x, "x"); // CHECK: x is present
- // FIXME: mapper should not map s.dummy or s.p; will be fixed when mapper
- // emits attach-style maps for pointer members.
- print_status(&s1.dummy, "dummy"); // CHECK: dummy is present
- // EXPECTED: dummy is not present
- // FIXME: mapper should not map s.dummy or s.p; will be fixed when mapper
- // emits attach-style maps for pointer members.
- print_status(&s1.p, "p"); // CHECK: p is present
- // EXPECTED: p is not present
- print_status(&s1.p[0], "p[0]"); // CHECK: p[0] 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");
#pragma omp target exit data map(delete : s1)
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 1b91d87071cbd..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
@@ -42,10 +42,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
- print_status(&s2.s1p->p, "p"); // CHECK: p is present
- // FIXME: mapper should emit attach-style maps for pointer members.
- // EXPECTED: p is not 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
index 09eed525d4340..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
@@ -6,6 +6,10 @@
// Array variant of mapper_map_ptee_only_2_ptr_indirections.c.
// 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];
@@ -45,11 +49,8 @@ int main() {
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
- // FIXME: mapper should not map s1p->p; will be fixed when mapper emits
- // attach-style maps for pointer members.
print_status(&s2arr[0].s1p->p,
- "s2arr[0].p"); // CHECK: s2arr[0].p is present
- // EXPECTED: s2arr[0].p is not present
+ "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
@@ -57,11 +58,8 @@ int main() {
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
- // FIXME: mapper should not map s1p->p; will be fixed when mapper emits
- // attach-style maps for pointer members.
print_status(&s2arr[1].s1p->p,
- "s2arr[1].p"); // CHECK: s2arr[1].p is present
- // EXPECTED: s2arr[1].p is not present
+ "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");
diff --git a/offload/test/mapping/mapper_map_ptee_only_2ndlevel.c b/offload/test/mapping/mapper_map_ptee_only_2ndlevel.c
index d7803b4e24d9a..2d5404a6d0cea 100644
--- a/offload/test/mapping/mapper_map_ptee_only_2ndlevel.c
+++ b/offload/test/mapping/mapper_map_ptee_only_2ndlevel.c
@@ -35,18 +35,14 @@ 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
- print_status(&s2.s1.p, "p"); // CHECK: p is present
- // FIXME: mapper should emit attach-style maps for pointer members.
- // EXPECTED: p is not 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)
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 present
- // FIXME: mapper should emit attach-style maps for pointer members.
- // EXPECTED: p is not present
- print_status(&s2.s1.p[0], "p[0]"); // CHECK: p[0] 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
}
>From 7ed335e622b66636490ea6d8cdd9292849dcae51 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Sun, 2 Aug 2026 23:10:14 -0700
Subject: [PATCH 2/2] [OpenMP][NFC] Document HasAttachPtr semantics in mapper
codegen comments
Address review comments:
- Add the struct declarations for the types used by the two worked
examples in emitUserDefinedMapper, so the entry tables can be read
without reconstructing the types from the entry sizes.
- Spell out which entries HasAttachPtr is set on: every entry whose
storage lies in a pointee block, i.e. the combined entry for the
block and the individual member entries that are MEMBER_OF it, but
never the ATTACH entry itself, and never an entry that maps the
pointer as an object in its own right.
Comment-only change.
Co-Authored-By: Claude Opus 5 <noreply at anthropic.com>
---
.../llvm/Frontend/OpenMP/OMPIRBuilder.h | 20 +++++++++++++++++++
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 18 ++++++++++++++---
2 files changed, 35 insertions(+), 3 deletions(-)
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
index 2fd98feea0c33..c4aaf74a227fd 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
@@ -2944,6 +2944,26 @@ class OpenMPIRBuilder {
MapNamesArrayTy Names;
/// True for entries that have an attach ptr, and thus an accompanying
/// ATTACH entry linking that ptr to its ptee.
+ ///
+ /// This is a property of the storage block an entry describes, not of the
+ /// map clause list item: it is true iff the entry's storage is the pointee
+ /// storage reached through some attach ptr. So, for `int *p; map(p[1:10])`,
+ /// which produces
+ /// &p[1], &p[1], 10*sizeof(int), TO|FROM <- HasAttachPtr = true
+ /// &p, &p[1], sizeof(void*), ATTACH <- HasAttachPtr = false
+ /// it is set on the pointee entry only. It is never set on the ATTACH entry
+ /// itself, nor on an entry that maps the pointer as an object in its own
+ /// right (e.g. the `map(p)` entry for the pointer's own storage).
+ ///
+ /// It is set on every entry whose storage lies in a pointee block,
+ /// including a combined struct entry for such a block and the individual
+ /// member entries that are MEMBER_OF it. e.g. for
+ /// `map(s2.s1p->x, s2.s1p->y)`:
+ /// &s2.s1p[0], &s2.s1p->x, sizeof(x..y), ALLOC <- true
+ /// &s2.s1p[0], &s2.s1p->x, 4, MEMBER_OF(1)|TO|FROM <- true
+ /// &s2.s1p[0], &s2.s1p->y, 4, MEMBER_OF(1)|TO|FROM <- true
+ /// &s2.s1p, &s2.s1p->x, sizeof(void*), ATTACH <- false
+ /// with s2.s1p as the attach ptr for all three.
MapHasAttachPtrArrayTy HasAttachPtr;
StructNonContiguousInfo NonContigInfo;
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 3df035f75746c..11cb609c45663 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -10536,6 +10536,8 @@ Expected<Function *> OpenMPIRBuilder::emitUserDefinedMapper(
// start).
//
// Example 1:
+ // struct S { int x; int *p; };
+ //
// mapper: #pragma omp declare mapper(id: S s) map(s.x, s.p[0:10])
// use: S arr[2]; ... map(arr)
// entries per element:
@@ -10545,6 +10547,9 @@ Expected<Function *> OpenMPIRBuilder::emitUserDefinedMapper(
// &arr[i].p, &arr[i].p[0], sizeof(int*), ATTACH (**)
//
// Example 2:
+ // struct S1 { int x; int y; };
+ // struct S2 { int z; S1 *s1p; };
+ //
// mapper: #pragma omp declare mapper(S2 s2) map(s2.z, s2.s1p->x,
// s2.s1p->y)
// use: S2 arr[2]; ... map(arr)
@@ -10552,19 +10557,26 @@ Expected<Function *> OpenMPIRBuilder::emitUserDefinedMapper(
//
// &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, 4, MEMBER_OF(N+2)|TO|FROM (***)
- // &arr[i].s1p[0], &arr[i].s1p->y, 4, MEMBER_OF(N+2)|TO|FROM (***)
+ // &arr[i].s1p[0], &arr[i].s1p->x, 4, MEMBER_OF(N+2)|TO|FROM (*)(***)
+ // &arr[i].s1p[0], &arr[i].s1p->y, 4, 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).
//
+ // HasAttachPtr is set on all of the s1p entries except the ATTACH one:
+ // the combined ALLOC entry for the s1p->x..y block, and the individual
+ // x/y entries that are MEMBER_OF that block, all describe storage
+ // reached through the attach ptr arr[i].s1p.
+ //
// Entries of the following kinds do NOT receive a new outer MEMBER_OF
// linking them to the parent struct:
//
// * (*) Entries with HasAttachPtr: they represent pointee data that
// occupies a different storage block than the struct being mapped, so
- // they are not a member of it.
+ // they are not a member of it. They may still be MEMBER_OF an entry
+ // within that pointee block, in which case those pre-existing bits are
+ // shifted -- see (***).
// * (**) ATTACH entries: they are not a member of anything — they just
// link a ptr to its ptee.
// * All entries when PreserveMemberOfFlags is set (the Flang/MLIR path):
More information about the llvm-commits
mailing list