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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
