Author: Abhinav Gaba Date: 2026-08-03T18:23:31-07:00 New Revision: 04a123b17c1ab097e3475da53a7e485c7b184a04
URL: https://github.com/llvm/llvm-project/commit/04a123b17c1ab097e3475da53a7e485c7b184a04 DIFF: https://github.com/llvm/llvm-project/commit/04a123b17c1ab097e3475da53a7e485c7b184a04.diff LOG: [OpenMP][Clang] Enable `ATTACH`-style maps for mappers. (#210213) 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. --------- Co-authored-by: Claude Opus 4.8 <[email protected]> Added: Modified: clang/docs/OpenMPSupport.md clang/docs/ReleaseNotes.md clang/lib/CodeGen/CGOpenMPRuntime.cpp clang/test/OpenMP/declare_mapper_codegen.cpp clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp offload/test/mapping/mapper_enter_data_always_present_ptee.c offload/test/mapping/mapper_map_mbr_ptee_then_present_mbr_ptee.c offload/test/mapping/mapper_map_mbr_then_present_mbr_ptee.c offload/test/mapping/mapper_map_present_ptee.c offload/test/mapping/mapper_map_ptee_only.c offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections.c offload/test/mapping/mapper_map_ptee_only_2_ptr_indirections_array.c offload/test/mapping/mapper_map_ptee_only_2ndlevel.c Removed: ################################################################################ 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 d5b5a014d9ef6..fc0aaeafc4fcf 100644 --- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h +++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h @@ -2960,6 +2960,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 @@ -2978,6 +2979,29 @@ 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. + /// + /// 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; /// Append arrays in \a CurInfo. @@ -2990,6 +3014,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 6fe6f268f1ef5..db55fc3ab0fa0 100644 --- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp @@ -10635,19 +10635,78 @@ 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: + // 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: + // + // &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: + // 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) + // 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). + // + // 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 diff erent storage block than the struct being mapped, so + // 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): + // 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; @@ -10758,12 +10817,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 90061026e6225..084dcb0a5847f 100644 --- a/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp +++ b/llvm/unittests/Frontend/OpenMPIRBuilderTest.cpp @@ -8345,6 +8345,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 dcaeec51c5b7a..07f46d53ad48c 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( @@ -8937,6 +8955,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 9da3d0afd326b..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: message: explicit extension not allowed - // CHECK: 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 0ee75ada5d00e..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: message: explicit extension not allowed - // CHECK: 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 ce6887a12f8de..7a0ed0b126c5a 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: message: device mapping required by 'present' motion modifier does not exist for host address - // CHECK: fatal error 1: failure of target construct while offloading is mandatory + // CHECK-OOB: message: device mapping required by 'present' motion modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes) + // CHECK-OOB: 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 } _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
