[clang] [llvm] [mlir] [OpenMP] Propagate PRESENT to pointee entries in mapper codegen (PR #210214)
Abhinav Gaba via llvm-commits
llvm-commits at lists.llvm.org
Tue Aug 4 12:06:06 PDT 2026
https://github.com/abhinavgaba updated https://github.com/llvm/llvm-project/pull/210214
>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/6] [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 9dfd09764d54e970c772fdbe17e19f112fb78301 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Fri, 24 Jul 2026 11:36:17 -0700
Subject: [PATCH 2/6] [OpenMP] Propagate PRESENT to pointee entries in mapper
codegen
Extend map-type-modifier propagation in emitUserDefinedMapper to the PRESENT
modifier, but only for entries that have an attach ptr (the pointee data, whose
storage differs from the struct's own). A present modifier on the outer clause
must require that pointee to be present on the device.
This is gated on a new PropagatePresentToPointee argument, which Clang sets from
CGM.getLangOpts().OpenMP >= 60. Before 6.0 the present modifier is treated as
not applying to the pointee: the spec committee confirmed the divergence
between the present motion modifier (to/from) and the present map-type modifier
(map) was unintentional, to be fixed as an OpenMP 6.0 erratum. Only propagation
is gated; present written directly in a mapper's own clause applies at all
versions.
A TODO notes PRESENT should also propagate to the struct's own members, which
is blocked while pointer members use PTR_AND_OBJ.
Update the present-check tests to their final 6.0-gated behavior.
Co-Authored-By: Claude Opus 4.8 <noreply at anthropic.com>
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 15 ++++--
...t_map_nested_ptr_member_mapper_codegen.cpp | 15 +++---
.../llvm/Frontend/OpenMP/OMPIRBuilder.h | 17 +++++--
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 47 ++++++++++++++-----
.../mapper_map_mbr_then_present_mbr_ptee.c | 27 ++++++-----
.../mapper_target_update_present_ptee.c | 27 ++++-------
6 files changed, 91 insertions(+), 57 deletions(-)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 6edabd4e42d88..3a4b5a351801a 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -10595,8 +10595,11 @@ getNestedDistributeDirective(ASTContext &Ctx, const OMPExecutableDirective &D) {
/// // 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);
+/// // bits have no meaning for ATTACH). PRESENT is additionally
+/// // propagated to pointee (attach-ptr) components at OpenMP >= 6.0.
+/// present_bit = (v60 && c.hasAttachPtr()) ? PRESENT : 0;
+/// imported_modifier_bits =
+/// type & (ALWAYS | DELETE | CLOSE | present_bit);
/// effective_type = c.isAttach() ? member_type
/// : member_type | imported_modifier_bits;
/// if (c.hasMapper())
@@ -10676,8 +10679,14 @@ void CGOpenMPRuntime::emitUserDefinedMapper(const OMPDeclareMapperDecl *D,
CGM.getCXXABI().getMangleContext().mangleCanonicalTypeName(Ty, Out);
std::string Name = getName({"omp_mapper", TyStr, D->getName()});
+ // Propagate the PRESENT modifier to pointee (attach-ptr) entries only for
+ // OpenMP >= 6.0; before 6.0 the present modifier does not apply to the
+ // pointee (see the OpenMP 6.0 erratum on the present motion vs. map-type
+ // modifier divergence).
+ bool PropagatePresentToPointee = CGM.getLangOpts().OpenMP >= 60;
llvm::Function *NewFn = cantFail(OMPBuilder.emitUserDefinedMapper(
- PrivatizeAndGenMapInfoCB, ElemTy, Name, CustomMapperCB));
+ PrivatizeAndGenMapInfoCB, ElemTy, Name, CustomMapperCB,
+ /*PreserveMemberOfFlags=*/false, PropagatePresentToPointee));
UDMMap.try_emplace(D, NewFn);
if (CGF)
FunctionUDMMap[CGF->CurFn].push_back(D);
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 821b5fd1652c4..40f22beb9bd7d 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
@@ -3,12 +3,9 @@
// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
-// PRESENT is propagated to pointee (attach-ptr) entries only at OpenMP >= 6.0.
-// FIXME: that propagation is done in a follow-on; until then the CHECK-60
-// output below is identical to CHECK (the pointee entries carry map-type mask
-// 1036 = ALWAYS|DELETE|CLOSE, without PRESENT). Once PRESENT is propagated, the
-// attach-ptr pointee entries should use mask 5132 = ALWAYS|DELETE|CLOSE|PRESENT
-// at 6.0.
+// PRESENT is propagated to pointee (attach-ptr) entries only at OpenMP >= 6.0:
+// at 6.0 those entries carry map-type mask 5132 = ALWAYS|DELETE|CLOSE|PRESENT,
+// while the default (<= 5.2) CHECK uses 1036 = ALWAYS|DELETE|CLOSE (no PRESENT).
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK-60
// expected-no-diagnostics
@@ -353,7 +350,7 @@ void foo(S2 *arr) {
// CHECK-60: br label [[OMP_TYPE_END9]]
// CHECK-60: omp.type.end9:
// 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: [[TMP38:%.*]] = and i64 [[TMP4]], 5132
// 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]]
@@ -377,7 +374,7 @@ void foo(S2 *arr) {
// CHECK-60: br label [[OMP_TYPE_END17]]
// CHECK-60: omp.type.end17:
// 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: [[TMP47:%.*]] = and i64 [[TMP4]], 5132
// 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]]
@@ -401,7 +398,7 @@ void foo(S2 *arr) {
// CHECK-60: br label [[OMP_TYPE_END25]]
// CHECK-60: omp.type.end25:
// 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: [[TMP56:%.*]] = and i64 [[TMP4]], 5132
// 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
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
index 2fd98feea0c33..3e46f74e71d74 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
@@ -3636,8 +3636,13 @@ class OpenMPIRBuilder {
/// // 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).
- /// imported_modifier_bits = type & (ALWAYS | DELETE | CLOSE);
+ /// // modifier bits have no meaning for ATTACH). PRESENT is
+ /// // additionally propagated to pointee (attach-ptr) components when
+ /// // PropagatePresentToPointee is set (OpenMP >= 6.0).
+ /// present_bit = (PropagatePresentToPointee && c.hasAttachPtr())
+ /// ? PRESENT : 0;
+ /// imported_modifier_bits = type & (ALWAYS | DELETE | CLOSE |
+ /// present_bit);
/// effective_type = c.isAttach() ? c.arg_type
/// : c.arg_type | imported_modifier_bits;
/// if (c.hasMapper())
@@ -3662,13 +3667,17 @@ class OpenMPIRBuilder {
/// \param FuncName Optional param to specify mapper function name.
/// \param CustomMapperCB Optional callback to generate code related to
/// custom mappers.
+ /// \param PropagatePresentToPointee If true, the PRESENT map-type modifier
+ /// from the outer clause is propagated to pointee (attach-ptr) entries the
+ /// mapper inserts. Callers set this only for OpenMP >= 6.0; at earlier
+ /// versions the present modifier is treated as not applying to the pointee.
LLVM_ABI Expected<Function *> emitUserDefinedMapper(
function_ref<MapInfosOrErrorTy(
InsertPointTy CodeGenIP, llvm::Value *PtrPHI, llvm::Value *BeginArg)>
PrivAndGenMapInfoCB,
llvm::Type *ElemTy, StringRef FuncName,
- CustomMapperCallbackTy CustomMapperCB,
- bool PreserveMemberOfFlags = false);
+ CustomMapperCallbackTy CustomMapperCB, bool PreserveMemberOfFlags = false,
+ bool PropagatePresentToPointee = false);
/// Generator for '#omp target data'
///
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 3df035f75746c..cb3cdcb601bc3 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -10426,7 +10426,7 @@ Expected<Function *> OpenMPIRBuilder::emitUserDefinedMapper(
llvm::Value *BeginArg)>
GenMapInfoCB,
Type *ElemTy, StringRef FuncName, CustomMapperCallbackTy CustomMapperCB,
- bool PreserveMemberOfFlags) {
+ bool PreserveMemberOfFlags, bool PropagatePresentToPointee) {
SmallVector<Type *> Params;
Params.emplace_back(Builder.getPtrTy());
Params.emplace_back(Builder.getPtrTy());
@@ -10673,16 +10673,41 @@ Expected<Function *> OpenMPIRBuilder::emitUserDefinedMapper(
// specified in the declared mapper.
//
// Map-type-modifying bits: ALWAYS, DELETE, CLOSE, PRESENT.
- // TODO: PRESENT is not propagated here yet. Doing so requires
- // distinguishing pointee entries from the struct's own storage; it is
- // handled in a follow-on.
- Value *ImportedModifierBits = Builder.CreateAnd(
- MapType,
- Builder.getInt64(
- static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>(
- OpenMPOffloadMappingFlags::OMP_MAP_ALWAYS |
- OpenMPOffloadMappingFlags::OMP_MAP_DELETE |
- OpenMPOffloadMappingFlags::OMP_MAP_CLOSE)));
+ //
+ // ALWAYS/DELETE/CLOSE are propagated to every (non-ATTACH) entry.
+ //
+ // PRESENT is propagated only to entries that have an attach ptr
+ // (HasAttachPtr): the pointee data, which occupies a different storage
+ // block than the struct being mapped and so is not covered by the
+ // present-check on the struct's own storage. A present modifier on the
+ // outer clause must still require that pointee to be present on the device.
+ //
+ // This is gated on \p PropagatePresentToPointee (set by callers only for
+ // OpenMP >= 6.0). Before 6.0 the present modifier is treated as not
+ // applying to the pointee: the spec committee confirmed the divergence
+ // between the present "motion" modifier (to/from) and the present map-type
+ // modifier (map) was unintentional, to be fixed as an OpenMP 6.0 erratum,
+ // so for 5.2 present is ignored for the pointee for both map and to/from.
+ //
+ // TODO: PRESENT should also be propagated to the struct's own members
+ // (e.g. the s.x, s.y of map(present, mapper(id): s)) so that an absent
+ // member triggers the present-check. We cannot do that yet: while pointer
+ // members are mapped with PTR_AND_OBJ, a single combined entry allocates
+ // the whole struct (including the pointer's storage), so propagating
+ // PRESENT to it would wrongly require the pointer's pointee to be present.
+ // Enable member propagation once Clang stops emitting PTR_AND_OBJ and uses
+ // attach-style maps throughout.
+ uint64_t ModifierBits =
+ static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>(
+ OpenMPOffloadMappingFlags::OMP_MAP_ALWAYS |
+ OpenMPOffloadMappingFlags::OMP_MAP_DELETE |
+ OpenMPOffloadMappingFlags::OMP_MAP_CLOSE);
+ if (PropagatePresentToPointee && Info->HasAttachPtr[I])
+ ModifierBits |=
+ static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>(
+ OpenMPOffloadMappingFlags::OMP_MAP_PRESENT);
+ Value *ImportedModifierBits =
+ Builder.CreateAnd(MapType, Builder.getInt64(ModifierBits));
Value *CurMapTypeWithModifiers = Builder.CreateOr(
CurMapType, ImportedModifierBits, "omp.maptype.with.modifiers");
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 bb8b6aa3c3e76..8d44e27146100 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,18 +1,16 @@
// 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
-// passes.
-//
-// 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].
+// present, so the propagated present modifier must fail the check -- but only
+// at OpenMP >= 6.0, since present is not propagated to the pointee before then.
+
+// OpenMP <= 5.2: present is not propagated to the pointee, so the run succeeds.
// RUN: %libomptarget-compile-generic -fopenmp-version=52
// RUN: %libomptarget-run-generic 2>&1 \
// RUN: | %fcheck-generic --check-prefixes=CHECK,CHECK-52
+
+// OpenMP 6.0: present is propagated to the pointee; the check fails.
// RUN: %libomptarget-compile-generic -fopenmp-version=60
-// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: %libomptarget-run-fail-generic 2>&1 \
// RUN: | %fcheck-generic --check-prefixes=CHECK,CHECK-60
#include <omp.h>
@@ -48,11 +46,14 @@ int main() {
print_status(&s1.p[0], "p[0]"); // CHECK: p[0] is not present
#pragma omp target enter data map(present, alloc : s1)
- // 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.
+ // 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
+ //
+ // clang-format off
+ // CHECK-60: omptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes)
+ // CHECK-60: omptarget fatal error 1: failure of target construct while offloading is mandatory
+ // clang-format on
fprintf(stderr, "done\n");
}
diff --git a/offload/test/mapping/mapper_target_update_present_ptee.c b/offload/test/mapping/mapper_target_update_present_ptee.c
index 10f11392ed76c..f241be882b6b7 100644
--- a/offload/test/mapping/mapper_target_update_present_ptee.c
+++ b/offload/test/mapping/mapper_target_update_present_ptee.c
@@ -6,17 +6,14 @@
// RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefix=CHECK-60
// Out-of-bounds: s.p[0:20] extends beyond the mapped region x[0:10]. At OpenMP
-// 6.0 the present modifier should be propagated to the pointee s.p[0:20], so
-// this should run-fail with a present error; at <= 5.2 present is (correctly)
-// not applied to the pointee, so the run succeeds.
-// FIXME: the present modifier is not yet propagated to the pointee, so the OOB
-// run currently succeeds at 6.0 too; propagating it is done in a follow-on.
-// EXPECTED (6.0): run-fail with a present-modifier error for s.p[0:20].
+// 6.0 the present modifier is propagated to the pointee s.p[0:20], so the
+// update fails the present check; at <= 5.2 present is not applied to the
+// pointee, so it succeeds.
// RUN: %libomptarget-compile-generic -fopenmp-version=52 -DOUT_OF_BOUNDS
// RUN: %libomptarget-run-generic 2>&1 \
// RUN: | %fcheck-generic --check-prefix=CHECK-52-OOB
// RUN: %libomptarget-compile-generic -fopenmp-version=60 -DOUT_OF_BOUNDS
-// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: %libomptarget-run-fail-generic 2>&1 \
// RUN: | %fcheck-generic --check-prefix=CHECK-60-OOB
#include <stdio.h>
@@ -54,6 +51,11 @@ int main() {
// CHECK-60-OOB: addr=0x[[#%x,HOST_ADDR:]], size=[[#%u,SIZE:]]
fprintf(stderr, "addr=%p, size=%zu\n", &s.p[0], 20 * sizeof(s.p[0]));
+ // At 6.0 the out-of-bounds pointee fails the present check inside f1().
+ // clang-format off
+ // CHECK-60-OOB: omptarget message: device mapping required by 'present' motion modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes)
+ // CHECK-60-OOB: omptarget fatal error 1: failure of target construct while offloading is mandatory
+ // clang-format on
#pragma omp target data map(from : s.y, x)
{
f1();
@@ -64,21 +66,12 @@ int main() {
// CHECK-60: 333 333
fprintf(stderr, "%d %d\n", x[0], s.y);
- // Out-of-bounds: present is not yet propagated to the pointee, so both
- // versions currently complete past the update.
- //
- // 5.2: present is (correctly) never applied to the pointee, so the update
+ // 5.2 out-of-bounds: present is not applied to the pointee, so the update
// completes.
// FIXME: even at 5.2, the update should still have happened for the subset of
// the pointee that is present (x[0:10]), instead of being silently ignored
// (x[0] currently reads back as garbage). Once that is fixed:
// EXPECTED-52-OOB: 333 333
// CHECK-52-OOB: done
- //
- // 6.0: present should be propagated to the pointee, so the update should fail
- // the present check for s.p[0:20]. That propagation is done in a follow-on;
- // until then the run completes.
- // EXPECTED-60-OOB: run-fail with a present-modifier error for s.p[0:20].
- // CHECK-60-OOB: done
fprintf(stderr, "done\n");
}
>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 3/6] [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):
>From 0dde5f20f5a9b88e19ead03723be60a1053215a6 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Mon, 3 Aug 2026 17:26:38 -0700
Subject: [PATCH 4/6] [OpenMP][NFC] Drop `omptarget` prefix from CHECK lines in
mapper present tests
Upstream 58f386207ac8 ("[offload] Remove `omptarget` references from
tests") made offload test CHECK lines generic so libomptarget components
can be moved/renamed -- some debug prints will come from `ompaccsupport`
rather than `omptarget`.
The two tests updated here re-add CHECK lines to files whose other
`omptarget`-prefixed lines that commit had already rewritten, so they
merged cleanly while reintroducing the old prefix. Match the convention
used by every other offload test. The address/size captures are
unchanged; only the component prefix is dropped.
---
offload/test/mapping/mapper_map_mbr_then_present_mbr_ptee.c | 4 ++--
offload/test/mapping/mapper_target_update_present_ptee.c | 4 ++--
2 files changed, 4 insertions(+), 4 deletions(-)
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 8d44e27146100..3e1856c2567ce 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
@@ -51,8 +51,8 @@ int main() {
// CHECK-52: done
//
// clang-format off
- // CHECK-60: omptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes)
- // CHECK-60: omptarget fatal error 1: failure of target construct while offloading is mandatory
+ // CHECK-60: message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes)
+ // CHECK-60: fatal error 1: failure of target construct while offloading is mandatory
// clang-format on
fprintf(stderr, "done\n");
diff --git a/offload/test/mapping/mapper_target_update_present_ptee.c b/offload/test/mapping/mapper_target_update_present_ptee.c
index f241be882b6b7..a3adcd6eb196a 100644
--- a/offload/test/mapping/mapper_target_update_present_ptee.c
+++ b/offload/test/mapping/mapper_target_update_present_ptee.c
@@ -53,8 +53,8 @@ int main() {
// At 6.0 the out-of-bounds pointee fails the present check inside f1().
// clang-format off
- // CHECK-60-OOB: omptarget message: device mapping required by 'present' motion modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes)
- // CHECK-60-OOB: omptarget fatal error 1: failure of target construct while offloading is mandatory
+ // CHECK-60-OOB: message: device mapping required by 'present' motion modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes)
+ // CHECK-60-OOB: fatal error 1: failure of target construct while offloading is mandatory
// clang-format on
#pragma omp target data map(from : s.y, x)
{
>From 98fa52b656d31b7d047009da0d095c0bd2f59c65 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Tue, 4 Aug 2026 00:16:49 -0700
Subject: [PATCH 5/6] [OpenMP][NFC] Align pointee-entry terminology with
HasAttachPtr docs
Drop the "pointee (attach-ptr) entries" phrasing in favor of the wording
already used by the landed HasAttachPtr documentation: name the flag
("entries with HasAttachPtr") where the comment sits next to a
HasAttachPtr/hasAttachPtr() test, and use plain "pointee entries" in
prose.
"attach-ptr entries" reads as the entries for the attach pointer itself,
which is exactly the case HasAttachPtr excludes -- it is false on the
ATTACH entry and on an entry mapping the pointer as an object in its own
right, and true only for the pointee storage reached through that ptr.
Comments only; no functional change.
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 11 ++++++-----
.../target_map_nested_ptr_member_mapper_codegen.cpp | 2 +-
llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h | 12 +++++++-----
3 files changed, 14 insertions(+), 11 deletions(-)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 3a4b5a351801a..51f1e8b22fcfc 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -10596,7 +10596,8 @@ getNestedDistributeDirective(ASTContext &Ctx, const OMPExecutableDirective &D) {
/// // 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 additionally
-/// // propagated to pointee (attach-ptr) components at OpenMP >= 6.0.
+/// // propagated to components with HasAttachPtr (the pointee data) at
+/// // OpenMP >= 6.0.
/// present_bit = (v60 && c.hasAttachPtr()) ? PRESENT : 0;
/// imported_modifier_bits =
/// type & (ALWAYS | DELETE | CLOSE | present_bit);
@@ -10679,10 +10680,10 @@ void CGOpenMPRuntime::emitUserDefinedMapper(const OMPDeclareMapperDecl *D,
CGM.getCXXABI().getMangleContext().mangleCanonicalTypeName(Ty, Out);
std::string Name = getName({"omp_mapper", TyStr, D->getName()});
- // Propagate the PRESENT modifier to pointee (attach-ptr) entries only for
- // OpenMP >= 6.0; before 6.0 the present modifier does not apply to the
- // pointee (see the OpenMP 6.0 erratum on the present motion vs. map-type
- // modifier divergence).
+ // Propagate the PRESENT modifier to the pointee entries (those with
+ // HasAttachPtr) only for OpenMP >= 6.0; before 6.0 the present modifier does
+ // not apply to the pointee (see the OpenMP 6.0 erratum on the present motion
+ // vs. map-type modifier divergence).
bool PropagatePresentToPointee = CGM.getLangOpts().OpenMP >= 60;
llvm::Function *NewFn = cantFail(OMPBuilder.emitUserDefinedMapper(
PrivatizeAndGenMapInfoCB, ElemTy, Name, CustomMapperCB,
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 40f22beb9bd7d..18d5133098162 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
@@ -3,7 +3,7 @@
// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
-// PRESENT is propagated to pointee (attach-ptr) entries only at OpenMP >= 6.0:
+// PRESENT is propagated to pointee entries only at OpenMP >= 6.0:
// at 6.0 those entries carry map-type mask 5132 = ALWAYS|DELETE|CLOSE|PRESENT,
// while the default (<= 5.2) CHECK uses 1036 = ALWAYS|DELETE|CLOSE (no PRESENT).
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK-60
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
index a892653dd1aa0..1965f7b983805 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
@@ -3694,8 +3694,9 @@ class OpenMPIRBuilder {
/// // 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
- /// // additionally propagated to pointee (attach-ptr) components when
- /// // PropagatePresentToPointee is set (OpenMP >= 6.0).
+ /// // additionally propagated to components with HasAttachPtr (the
+ /// // pointee data) when PropagatePresentToPointee is set
+ /// // (OpenMP >= 6.0).
/// present_bit = (PropagatePresentToPointee && c.hasAttachPtr())
/// ? PRESENT : 0;
/// imported_modifier_bits = type & (ALWAYS | DELETE | CLOSE |
@@ -3725,9 +3726,10 @@ class OpenMPIRBuilder {
/// \param CustomMapperCB Optional callback to generate code related to
/// custom mappers.
/// \param PropagatePresentToPointee If true, the PRESENT map-type modifier
- /// from the outer clause is propagated to pointee (attach-ptr) entries the
- /// mapper inserts. Callers set this only for OpenMP >= 6.0; at earlier
- /// versions the present modifier is treated as not applying to the pointee.
+ /// from the outer clause is propagated to the pointee entries the mapper
+ /// inserts, i.e. those with HasAttachPtr. Callers set this only for
+ /// OpenMP >= 6.0; at earlier versions the present modifier is treated as not
+ /// applying to the pointee.
LLVM_ABI Expected<Function *> emitUserDefinedMapper(
function_ref<MapInfosOrErrorTy(
InsertPointTy CodeGenIP, llvm::Value *PtrPHI, llvm::Value *BeginArg)>
>From e96646304dd8c81ec086194b429e17dae28ba46e Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Tue, 4 Aug 2026 12:04:27 -0700
Subject: [PATCH 6/6] [OpenMP][NFC] Refresh stale entry-layout comments in
mapper codegen test
The per-element entry list described the pre-attach-style codegen: it
listed four entries instead of five, ordered the combined entry before
the s2.z one, used MEMBER_OF(N+1) where MEMBER_OF(N+2) is emitted, and
marked the s1p->x/y entries PTR_AND_OBJ. It also carried a FIXME asking
for the attach-style codegen that is now in place.
Update it to the entries actually emitted, verified against the IR and
the test's own CHECK lines, and drop the obsolete FIXME. Also fix the
comment in foo(), which listed a single top-level entry when there are
two (the mapped array plus its ATTACH entry, per .offload_maptypes).
Comments only; no functional change.
---
...t_map_nested_ptr_member_mapper_codegen.cpp | 22 ++++++++++++-------
1 file changed, 14 insertions(+), 8 deletions(-)
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 18d5133098162..7e3a830ec2323 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
@@ -14,12 +14,17 @@
// S2 mapper for: map(to: arr[0:2])
// Per-element entries (i = array element index; N = __tgt_mapper_num_components()):
-// &arr[i], &arr[i].s1p, sizeof(s1p..z), MEMBER_OF(N) | ALLOC
-// &arr[i], &arr[i].z, sizeof(int), MEMBER_OF(N+1) | TO | FROM
-// &arr[i].s1p, &arr[i].s1p->x, sizeof(int), MEMBER_OF(N+1) | TO | FROM | PTR_AND_OBJ
-// &arr[i].s1p, &arr[i].s1p->y, sizeof(int), MEMBER_OF(N+1) | TO | FROM | PTR_AND_OBJ
-// FIXME: should use attach-style codegen for s1p->x/y instead of PTR_AND_OBJ,
-// which unnecessarily also maps s1p.
+// &arr[i], &arr[i].z, sizeof(int), MEMBER_OF(N)|TO|FROM
+// &arr[i].s1p[0], &arr[i].s1p->x, sizeof(s1p->x..y), ALLOC
+// &arr[i].s1p[0], &arr[i].s1p->x, sizeof(int), MEMBER_OF(N+2)|TO|FROM
+// &arr[i].s1p[0], &arr[i].s1p->y, sizeof(int), MEMBER_OF(N+2)|TO|FROM
+// &arr[i].s1p, &arr[i].s1p->x, sizeof(S1 *), ATTACH
+//
+// The s1p->x/y entries use attach-style codegen: their storage is the pointee
+// block, so they are not MEMBER_OF the arr[i] struct. They are MEMBER_OF the
+// combined ALLOC entry that covers that block (inner MEMBER_OF(2), shifted by
+// N), and a separate ATTACH entry links arr[i].s1p to it. Unlike PTR_AND_OBJ,
+// this does not also map s1p itself.
typedef struct {
int x;
@@ -34,8 +39,9 @@ typedef struct {
#pragma omp declare mapper(default : S2 s2) map(s2.z, s2.s1p->x, s2.s1p->y)
void foo(S2 *arr) {
- // &arr, &arr[0], 2*sizeof(S2), TARGET_PARAM | TO
- // (mapper handles individual members)
+ // Top-level entries (the mapper above expands each element's members):
+ // arr, &arr[0], 2*sizeof(S2), TO (with the S2 mapper attached)
+ // &arr, &arr[0], sizeof(S2 *), ATTACH (no mapper)
#pragma omp target enter data map(to: arr[0:2])
{}
}
More information about the llvm-commits
mailing list