[clang] [llvm] [mlir] [OpenMP][Clang] Enable `ATTACH`-style maps for mappers. (PR #210213)
via cfe-commits
cfe-commits at lists.llvm.org
Mon Aug 3 01:23:10 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-mlir-llvm
Author: Abhinav Gaba (abhinavgaba)
<details>
<summary>Changes</summary>
This is a follow-up to #<!-- -->153683 to support OpenMP compliant pointer-attachment
in `declare_mappers` via `ATTACH`-style maps.
In addition to enabling attach-style maps, we also need to propagate information about
which map entries are for "pointee" data, i.e. have an "attach-ptr", and thus occupy a different storage block than the base variable for which the mapper is being generated. e.g.
```c
S sa[10];
#pragma omp declare_mapper (default: S s) map (s.x, s.p[0:10])
#pragma omp target_enter_data map(sa)
```
The entry emitted for `s.p[0:10]` is for the pointee, i.e. it does not share storage with `s`.
Mapper codegen needs to know that the entry for `sa[1].p[0:10]`, for example, is not a `MEMBER_OF` the map of `sa`, as it occupies its own storage and has its own ref-count tracking etc.
Flang currently passes an unconditional `PreserveMemberofFlags` bool to the OMPIRBuilder function, which should eventually be propagated to using the per-entry information so that `map(s%x)` should get MEMBER_OF during mapper codgen, but `map(s%p(1:10))` should not. Currently, the per-entry MapInfo field is set to `false` for flang, so the change is a no-op for it. I don't have enough flang expertise/testing resources, so I'll let Andrew update Flang in follow-up changes.
---
Patch is 96.62 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/210213.diff
17 Files Affected:
- (modified) clang/docs/OpenMPSupport.md (+3-2)
- (modified) clang/docs/ReleaseNotes.md (+3)
- (modified) clang/lib/CodeGen/CGOpenMPRuntime.cpp (+65-5)
- (modified) clang/test/OpenMP/declare_mapper_codegen.cpp (+68-60)
- (modified) clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp (+196-160)
- (modified) llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h (+26)
- (modified) llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp (+68-15)
- (modified) llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp (+1)
- (modified) mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp (+20)
- (modified) offload/test/mapping/mapper_enter_data_always_present_ptee.c (+18-24)
- (modified) offload/test/mapping/mapper_map_mbr_ptee_then_present_mbr_ptee.c (+15-27)
- (modified) offload/test/mapping/mapper_map_mbr_then_present_mbr_ptee.c (+18-22)
- (modified) offload/test/mapping/mapper_map_present_ptee.c (+9-17)
- (modified) offload/test/mapping/mapper_map_ptee_only.c (+4-10)
- (modified) offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections.c (+2-4)
- (modified) offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections_array.c (+6-8)
- (modified) offload/test/mapping/mapper_map_ptee_only_2ndlevel.c (+4-8)
``````````diff
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_ma...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/210213
More information about the cfe-commits
mailing list