diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -7157,6 +7157,7 @@ 0, Address::invalid()}; Address Base = Address::invalid(); bool IsArraySection = false; + bool HasCompleteRecord = false; }; private: @@ -7339,11 +7340,10 @@ if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_close) != MapModifiers.end()) Bits |= OMP_MAP_CLOSE; - if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_present) - != MapModifiers.end()) - Bits |= OMP_MAP_PRESENT; - if (llvm::find(MotionModifiers, OMPC_MOTION_MODIFIER_present) - != MotionModifiers.end()) + if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_present) != + MapModifiers.end() || + llvm::find(MotionModifiers, OMPC_MOTION_MODIFIER_present) != + MotionModifiers.end()) Bits |= OMP_MAP_PRESENT; if (IsNonContiguous) Bits |= OMP_MAP_NON_CONTIG; @@ -7921,6 +7921,10 @@ FirstPointerInComplexData = false; } } + // If ran into the whole component - allocate the space for the whole + // record. + if (!EncounteredME) + PartialStruct.HasCompleteRecord = true; if (!IsNonContiguous) return; @@ -8205,151 +8209,100 @@ } } -public: - MappableExprsHandler(const OMPExecutableDirective &Dir, CodeGenFunction &CGF) - : CurDir(&Dir), CGF(CGF) { - // Extract firstprivate clause information. - for (const auto *C : Dir.getClausesOfKind()) - for (const auto *D : C->varlists()) - FirstPrivateDecls.try_emplace( - cast(cast(D)->getDecl()), C->isImplicit()); - // Extract implicit firstprivates from uses_allocators clauses. - for (const auto *C : Dir.getClausesOfKind()) { - for (unsigned I = 0, E = C->getNumberOfAllocators(); I < E; ++I) { - OMPUsesAllocatorsClause::Data D = C->getAllocatorData(I); - if (const auto *DRE = dyn_cast_or_null(D.AllocatorTraits)) - FirstPrivateDecls.try_emplace(cast(DRE->getDecl()), - /*Implicit=*/true); - else if (const auto *VD = dyn_cast( - cast(D.Allocator->IgnoreParenImpCasts()) - ->getDecl())) - FirstPrivateDecls.try_emplace(VD, /*Implicit=*/true); - } - } - // Extract device pointer clause information. - for (const auto *C : Dir.getClausesOfKind()) - for (auto L : C->component_lists()) - DevPointersMap[std::get<0>(L)].push_back(std::get<1>(L)); - } - - /// Constructor for the declare mapper directive. - MappableExprsHandler(const OMPDeclareMapperDecl &Dir, CodeGenFunction &CGF) - : CurDir(&Dir), CGF(CGF) {} - - /// 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 - /// individual struct members. - void emitCombinedEntry(MapCombinedInfoTy &CombinedInfo, - MapFlagsArrayTy &CurTypes, - const StructRangeInfoTy &PartialStruct, - const ValueDecl *VD = nullptr, - bool NotTargetParams = true) const { - if (CurTypes.size() == 1 && - ((CurTypes.back() & OMP_MAP_MEMBER_OF) != OMP_MAP_MEMBER_OF) && - !PartialStruct.IsArraySection) - return; - CombinedInfo.Exprs.push_back(VD); - // Base is the base of the struct - CombinedInfo.BasePointers.push_back(PartialStruct.Base.getPointer()); - // Pointer is the address of the lowest element - llvm::Value *LB = PartialStruct.LowestElem.second.getPointer(); - CombinedInfo.Pointers.push_back(LB); - // There should not be a mapper for a combined entry. - CombinedInfo.Mappers.push_back(nullptr); - // Size is (addr of {highest+1} element) - (addr of lowest element) - llvm::Value *HB = PartialStruct.HighestElem.second.getPointer(); - llvm::Value *HAddr = CGF.Builder.CreateConstGEP1_32(HB, /*Idx0=*/1); - llvm::Value *CLAddr = CGF.Builder.CreatePointerCast(LB, CGF.VoidPtrTy); - llvm::Value *CHAddr = CGF.Builder.CreatePointerCast(HAddr, CGF.VoidPtrTy); - llvm::Value *Diff = CGF.Builder.CreatePtrDiff(CHAddr, CLAddr); - llvm::Value *Size = CGF.Builder.CreateIntCast(Diff, CGF.Int64Ty, - /*isSigned=*/false); - CombinedInfo.Sizes.push_back(Size); - // Map type is always TARGET_PARAM, if generate info for captures. - CombinedInfo.Types.push_back(NotTargetParams ? OMP_MAP_NONE - : OMP_MAP_TARGET_PARAM); - // If any element has the present modifier, then make sure the runtime - // doesn't attempt to allocate the struct. - if (CurTypes.end() != - llvm::find_if(CurTypes, [](OpenMPOffloadMappingFlags Type) { - return Type & OMP_MAP_PRESENT; - })) - CombinedInfo.Types.back() |= OMP_MAP_PRESENT; - // Remove TARGET_PARAM flag from the first element if any. - if (!CurTypes.empty()) - CurTypes.front() &= ~OMP_MAP_TARGET_PARAM; - - // All other current entries will be MEMBER_OF the combined entry - // (except for PTR_AND_OBJ entries which do not have a placeholder value - // 0xFFFF in the MEMBER_OF field). - OpenMPOffloadMappingFlags MemberOfFlag = - getMemberOfFlag(CombinedInfo.BasePointers.size() - 1); - for (auto &M : CurTypes) - setCorrectMemberOfFlag(M, MemberOfFlag); - } - /// Generate all the base pointers, section pointers, sizes, map types, and /// mappers for the extracted mappable expressions (all included in \a /// CombinedInfo). Also, for each item that relates with a device pointer, a /// pair of the relevant declaration and index where it occurs is appended to /// the device pointers info array. - void generateAllInfo( - MapCombinedInfoTy &CombinedInfo, + void generateAllInfoForClauses( + ArrayRef Clauses, MapCombinedInfoTy &CombinedInfo, const llvm::DenseSet> &SkipVarSet = llvm::DenseSet>()) const { // We have to process the component lists that relate with the same // declaration in a single chunk so that we can generate the map flags // correctly. Therefore, we organize all lists in a map. - llvm::MapVector> Info; + enum MapKind { Present, Allocs, Other, Total }; + llvm::MapVector, + SmallVector, 4>> + Info; // Helper function to fill the information map for the different supported // clauses. auto &&InfoGen = [&Info, &SkipVarSet]( - const ValueDecl *D, + const ValueDecl *D, MapKind Kind, OMPClauseMappableExprCommon::MappableExprComponentListRef L, OpenMPMapClauseKind MapType, ArrayRef MapModifiers, ArrayRef MotionModifiers, bool ReturnDevicePointer, bool IsImplicit, const ValueDecl *Mapper, const Expr *VarRef = nullptr, bool ForDeviceAddr = false) { - const ValueDecl *VD = - D ? cast(D->getCanonicalDecl()) : nullptr; - if (SkipVarSet.count(VD)) + if (SkipVarSet.contains(D)) return; - Info[VD].emplace_back(L, MapType, MapModifiers, MotionModifiers, - ReturnDevicePointer, IsImplicit, Mapper, VarRef, - ForDeviceAddr); + auto It = Info.find(D); + if (It == Info.end()) + It = Info + .insert(std::make_pair( + D, SmallVector, 4>(Total))) + .first; + It->second[Kind].emplace_back( + L, MapType, MapModifiers, MotionModifiers, ReturnDevicePointer, + IsImplicit, Mapper, VarRef, ForDeviceAddr); }; - assert(CurDir.is() && - "Expect a executable directive"); - const auto *CurExecDir = CurDir.get(); - for (const auto *C : CurExecDir->getClausesOfKind()) { + for (const auto *Cl : Clauses) { + const auto *C = dyn_cast(Cl); + if (!C) + continue; + MapKind Kind = Other; + if (!C->getMapTypeModifiers().empty() && + llvm::any_of(C->getMapTypeModifiers(), [](OpenMPMapModifierKind K) { + return K == OMPC_MAP_MODIFIER_present; + })) + Kind = Present; + else if (C->getMapType() == OMPC_MAP_alloc) + Kind = Allocs; const auto *EI = C->getVarRefs().begin(); for (const auto L : C->component_lists()) { - // The Expression is not correct if the mapping is implicit const Expr *E = (C->getMapLoc().isValid()) ? *EI : nullptr; - InfoGen(std::get<0>(L), std::get<1>(L), C->getMapType(), + InfoGen(std::get<0>(L), Kind, std::get<1>(L), C->getMapType(), C->getMapTypeModifiers(), llvm::None, /*ReturnDevicePointer=*/false, C->isImplicit(), std::get<2>(L), E); ++EI; } } - for (const auto *C : CurExecDir->getClausesOfKind()) { + for (const auto *Cl : Clauses) { + const auto *C = dyn_cast(Cl); + if (!C) + continue; + MapKind Kind = Other; + if (!C->getMotionModifiers().empty() && + llvm::any_of(C->getMotionModifiers(), [](OpenMPMotionModifierKind K) { + return K == OMPC_MOTION_MODIFIER_present; + })) + Kind = Present; const auto *EI = C->getVarRefs().begin(); for (const auto L : C->component_lists()) { - InfoGen(std::get<0>(L), std::get<1>(L), OMPC_MAP_to, llvm::None, + InfoGen(std::get<0>(L), Kind, std::get<1>(L), OMPC_MAP_to, llvm::None, C->getMotionModifiers(), /*ReturnDevicePointer=*/false, C->isImplicit(), std::get<2>(L), *EI); ++EI; } } - for (const auto *C : CurExecDir->getClausesOfKind()) { + for (const auto *Cl : Clauses) { + const auto *C = dyn_cast(Cl); + if (!C) + continue; + MapKind Kind = Other; + if (!C->getMotionModifiers().empty() && + llvm::any_of(C->getMotionModifiers(), [](OpenMPMotionModifierKind K) { + return K == OMPC_MOTION_MODIFIER_present; + })) + Kind = Present; const auto *EI = C->getVarRefs().begin(); for (const auto L : C->component_lists()) { - InfoGen(std::get<0>(L), std::get<1>(L), OMPC_MAP_from, llvm::None, + InfoGen(std::get<0>(L), Kind, std::get<1>(L), OMPC_MAP_from, llvm::None, C->getMotionModifiers(), /*ReturnDevicePointer=*/false, C->isImplicit(), std::get<2>(L), *EI); ++EI; @@ -8362,12 +8315,15 @@ // section. It is the user fault if that was not mapped before. If there is // no map information and the pointer is a struct member, then we defer the // emission of that entry until the whole struct has been processed. - llvm::MapVector> + llvm::MapVector, + SmallVector> DeferredInfo; MapCombinedInfoTy UseDevicePtrCombinedInfo; - for (const auto *C : - CurExecDir->getClausesOfKind()) { + for (const auto *Cl : Clauses) { + const auto *C = dyn_cast(Cl); + if (!C) + continue; for (const auto L : C->component_lists()) { OMPClauseMappableExprCommon::MappableExprComponentListRef Components = std::get<1>(L); @@ -8384,28 +8340,34 @@ // We potentially have map information for this declaration already. // Look for the first set of components that refer to it. if (It != Info.end()) { - auto *CI = llvm::find_if(It->second, [VD](const MapInfo &MI) { - return MI.Components.back().getAssociatedDeclaration() == VD; - }); - // If we found a map entry, signal that the pointer has to be returned - // and move on to the next declaration. - // Exclude cases where the base pointer is mapped as array subscript, - // array section or array shaping. The base address is passed as a - // pointer to base in this case and cannot be used as a base for - // use_device_ptr list item. - if (CI != It->second.end()) { - auto PrevCI = std::next(CI->Components.rbegin()); - const auto *VarD = dyn_cast(VD); - if (CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory() || - isa(IE) || - !VD->getType().getNonReferenceType()->isPointerType() || - PrevCI == CI->Components.rend() || - isa(PrevCI->getAssociatedExpression()) || !VarD || - VarD->hasLocalStorage()) { - CI->ReturnDevicePointer = true; - continue; + bool Found = false; + for (auto &Data : It->second) { + auto *CI = llvm::find_if(Data, [VD](const MapInfo &MI) { + return MI.Components.back().getAssociatedDeclaration() == VD; + }); + // If we found a map entry, signal that the pointer has to be + // returned and move on to the next declaration. Exclude cases where + // the base pointer is mapped as array subscript, array section or + // array shaping. The base address is passed as a pointer to base in + // this case and cannot be used as a base for use_device_ptr list + // item. + if (CI != Data.end()) { + auto PrevCI = std::next(CI->Components.rbegin()); + const auto *VarD = dyn_cast(VD); + if (CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory() || + isa(IE) || + !VD->getType().getNonReferenceType()->isPointerType() || + PrevCI == CI->Components.rend() || + isa(PrevCI->getAssociatedExpression()) || !VarD || + VarD->hasLocalStorage()) { + CI->ReturnDevicePointer = true; + Found = true; + break; + } } } + if (Found) + continue; } // We didn't find any match in our map information - generate a zero @@ -8419,8 +8381,9 @@ // Nonetheless, generateInfoForComponentList must be called to take // the pointer into account for the calculation of the range of the // partial struct. - InfoGen(nullptr, Components, OMPC_MAP_unknown, llvm::None, llvm::None, - /*ReturnDevicePointer=*/false, C->isImplicit(), nullptr); + InfoGen(nullptr, Other, Components, OMPC_MAP_unknown, llvm::None, + llvm::None, /*ReturnDevicePointer=*/false, C->isImplicit(), + nullptr); DeferredInfo[nullptr].emplace_back(IE, VD, /*ForDeviceAddr=*/false); } else { llvm::Value *Ptr = @@ -8443,8 +8406,10 @@ // no map information and the pointer is a struct member, then we defer the // emission of that entry until the whole struct has been processed. llvm::SmallDenseSet, 4> Processed; - for (const auto *C : - CurExecDir->getClausesOfKind()) { + for (const auto *Cl : Clauses) { + const auto *C = dyn_cast(Cl); + if (!C) + continue; for (const auto L : C->component_lists()) { assert(!std::get<1>(L).empty() && "Not expecting empty list of components!"); @@ -8461,15 +8426,21 @@ // We potentially have map information for this declaration already. // Look for the first set of components that refer to it. if (It != Info.end()) { - auto *CI = llvm::find_if(It->second, [VD](const MapInfo &MI) { - return MI.Components.back().getAssociatedDeclaration() == VD; - }); - // If we found a map entry, signal that the pointer has to be returned - // and move on to the next declaration. - if (CI != It->second.end()) { - CI->ReturnDevicePointer = true; - continue; + bool Found = false; + for (auto &Data : It->second) { + auto *CI = llvm::find_if(Data, [VD](const MapInfo &MI) { + return MI.Components.back().getAssociatedDeclaration() == VD; + }); + // If we found a map entry, signal that the pointer has to be + // returned and move on to the next declaration. + if (CI != Data.end()) { + CI->ReturnDevicePointer = true; + Found = true; + break; + } } + if (Found) + continue; } // We didn't find any match in our map information - generate a zero @@ -8483,7 +8454,7 @@ // Nonetheless, generateInfoForComponentList must be called to take // the pointer into account for the calculation of the range of the // partial struct. - InfoGen(nullptr, std::get<1>(L), OMPC_MAP_unknown, llvm::None, + InfoGen(nullptr, Other, std::get<1>(L), OMPC_MAP_unknown, llvm::None, llvm::None, /*ReturnDevicePointer=*/false, C->isImplicit(), nullptr, nullptr, /*ForDeviceAddr=*/true); DeferredInfo[nullptr].emplace_back(IE, VD, /*ForDeviceAddr=*/true); @@ -8504,47 +8475,47 @@ } } - for (const auto &M : Info) { - // Underlying variable declaration used in the map clause. - const ValueDecl *VD = std::get<0>(M); - + for (const auto &Data : Info) { + StructRangeInfoTy PartialStruct; // Temporary generated information. MapCombinedInfoTy CurInfo; - StructRangeInfoTy PartialStruct; - - for (const MapInfo &L : M.second) { - assert(!L.Components.empty() && - "Not expecting declaration with no component lists."); - - // Remember the current base pointer index. - unsigned CurrentBasePointersIdx = CurInfo.BasePointers.size(); - CurInfo.NonContigInfo.IsNonContiguous = - L.Components.back().isNonContiguous(); - generateInfoForComponentList( - L.MapType, L.MapModifiers, L.MotionModifiers, L.Components, CurInfo, - PartialStruct, /*IsFirstComponentList=*/false, L.IsImplicit, - L.Mapper, L.ForDeviceAddr, VD, L.VarRef); - - // If this entry relates with a device pointer, set the relevant - // declaration and add the 'return pointer' flag. - if (L.ReturnDevicePointer) { - assert(CurInfo.BasePointers.size() > CurrentBasePointersIdx && - "Unexpected number of mapped base pointers."); - - const ValueDecl *RelevantVD = - L.Components.back().getAssociatedDeclaration(); - assert(RelevantVD && - "No relevant declaration related with device pointer??"); - - CurInfo.BasePointers[CurrentBasePointersIdx].setDevicePtrDecl( - RelevantVD); - CurInfo.Types[CurrentBasePointersIdx] |= OMP_MAP_RETURN_PARAM; + const Decl *D = Data.first; + const ValueDecl *VD = cast_or_null(D); + for (const auto &M : Data.second) { + for (const MapInfo &L : M) { + assert(!L.Components.empty() && + "Not expecting declaration with no component lists."); + + // Remember the current base pointer index. + unsigned CurrentBasePointersIdx = CurInfo.BasePointers.size(); + CurInfo.NonContigInfo.IsNonContiguous = + L.Components.back().isNonContiguous(); + generateInfoForComponentList( + L.MapType, L.MapModifiers, L.MotionModifiers, L.Components, + CurInfo, PartialStruct, /*IsFirstComponentList=*/false, + L.IsImplicit, L.Mapper, L.ForDeviceAddr, VD, L.VarRef); + + // If this entry relates with a device pointer, set the relevant + // declaration and add the 'return pointer' flag. + if (L.ReturnDevicePointer) { + assert(CurInfo.BasePointers.size() > CurrentBasePointersIdx && + "Unexpected number of mapped base pointers."); + + const ValueDecl *RelevantVD = + L.Components.back().getAssociatedDeclaration(); + assert(RelevantVD && + "No relevant declaration related with device pointer??"); + + CurInfo.BasePointers[CurrentBasePointersIdx].setDevicePtrDecl( + RelevantVD); + CurInfo.Types[CurrentBasePointersIdx] |= OMP_MAP_RETURN_PARAM; + } } } // Append any pending zero-length pointers which are struct members and // used with use_device_ptr or use_device_addr. - auto CI = DeferredInfo.find(M.first); + auto CI = DeferredInfo.find(Data.first); if (CI != DeferredInfo.end()) { for (const DeferredDevicePtrEntryTy &L : CI->second) { llvm::Value *BasePtr; @@ -8563,9 +8534,9 @@ BasePtr = this->CGF.EmitLValue(L.IE).getPointer(CGF); Ptr = this->CGF.EmitLoadOfScalar(this->CGF.EmitLValue(L.IE), L.IE->getExprLoc()); - // Entry is PTR_AND_OBJ and RETURN_PARAM. Also, set the placeholder - // value MEMBER_OF=FFFF so that the entry is later updated with the - // correct value of MEMBER_OF. + // Entry is PTR_AND_OBJ and RETURN_PARAM. Also, set the + // placeholder value MEMBER_OF=FFFF so that the entry is later + // updated with the correct value of MEMBER_OF. CurInfo.Types.push_back(OMP_MAP_PTR_AND_OBJ | OMP_MAP_RETURN_PARAM | OMP_MAP_MEMBER_OF); } @@ -8577,82 +8548,132 @@ CurInfo.Mappers.push_back(nullptr); } } - // If there is an entry in PartialStruct it means we have a struct with // individual members mapped. Emit an extra combined entry. - if (PartialStruct.Base.isValid()) + if (PartialStruct.Base.isValid()) { + CurInfo.NonContigInfo.Dims.push_back(0); emitCombinedEntry(CombinedInfo, CurInfo.Types, PartialStruct, VD); + } - // We need to append the results of this capture to what we already have. + // We need to append the results of this capture to what we already + // have. CombinedInfo.append(CurInfo); } // Append data for use_device_ptr clauses. CombinedInfo.append(UseDevicePtrCombinedInfo); } - /// Generate all the base pointers, section pointers, sizes, map types, and - /// mappers for the extracted map clauses of user-defined mapper (all included - /// in \a CombinedInfo). - void generateAllInfoForMapper(MapCombinedInfoTy &CombinedInfo) const { - assert(CurDir.is() && - "Expect a declare mapper directive"); - const auto *CurMapperDir = CurDir.get(); - // We have to process the component lists that relate with the same - // declaration in a single chunk so that we can generate the map flags - // correctly. Therefore, we organize all lists in a map. - llvm::MapVector> Info; - - // Fill the information map for map clauses. - for (const auto *C : CurMapperDir->clauselists()) { - const auto *MC = cast(C); - const auto *EI = MC->getVarRefs().begin(); - for (const auto L : MC->component_lists()) { - // The Expression is not correct if the mapping is implicit - const Expr *E = (MC->getMapLoc().isValid()) ? *EI : nullptr; - const ValueDecl *VD = - std::get<0>(L) ? cast(std::get<0>(L)->getCanonicalDecl()) - : nullptr; - // Get the corresponding user-defined mapper. - Info[VD].emplace_back(std::get<1>(L), MC->getMapType(), - MC->getMapTypeModifiers(), llvm::None, - /*ReturnDevicePointer=*/false, MC->isImplicit(), - std::get<2>(L), E); - ++EI; +public: + MappableExprsHandler(const OMPExecutableDirective &Dir, CodeGenFunction &CGF) + : CurDir(&Dir), CGF(CGF) { + // Extract firstprivate clause information. + for (const auto *C : Dir.getClausesOfKind()) + for (const auto *D : C->varlists()) + FirstPrivateDecls.try_emplace( + cast(cast(D)->getDecl()), C->isImplicit()); + // Extract implicit firstprivates from uses_allocators clauses. + for (const auto *C : Dir.getClausesOfKind()) { + for (unsigned I = 0, E = C->getNumberOfAllocators(); I < E; ++I) { + OMPUsesAllocatorsClause::Data D = C->getAllocatorData(I); + if (const auto *DRE = dyn_cast_or_null(D.AllocatorTraits)) + FirstPrivateDecls.try_emplace(cast(DRE->getDecl()), + /*Implicit=*/true); + else if (const auto *VD = dyn_cast( + cast(D.Allocator->IgnoreParenImpCasts()) + ->getDecl())) + FirstPrivateDecls.try_emplace(VD, /*Implicit=*/true); } } + // Extract device pointer clause information. + for (const auto *C : Dir.getClausesOfKind()) + for (auto L : C->component_lists()) + DevPointersMap[std::get<0>(L)].push_back(std::get<1>(L)); + } - for (const auto &M : Info) { - // We need to know when we generate information for the first component - // associated with a capture, because the mapping flags depend on it. - bool IsFirstComponentList = true; - - // Underlying variable declaration used in the map clause. - const ValueDecl *VD = std::get<0>(M); + /// Constructor for the declare mapper directive. + MappableExprsHandler(const OMPDeclareMapperDecl &Dir, CodeGenFunction &CGF) + : CurDir(&Dir), CGF(CGF) {} - // Temporary generated information. - MapCombinedInfoTy CurInfo; - StructRangeInfoTy PartialStruct; + /// 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 + /// individual struct members. + void emitCombinedEntry(MapCombinedInfoTy &CombinedInfo, + MapFlagsArrayTy &CurTypes, + const StructRangeInfoTy &PartialStruct, + const ValueDecl *VD = nullptr, + bool NotTargetParams = true) const { + if (CurTypes.size() == 1 && + ((CurTypes.back() & OMP_MAP_MEMBER_OF) != OMP_MAP_MEMBER_OF) && + !PartialStruct.IsArraySection) + return; + Address LBAddr = PartialStruct.LowestElem.second; + Address HBAddr = PartialStruct.HighestElem.second; + if (PartialStruct.HasCompleteRecord) { + LBAddr = PartialStruct.Base; + HBAddr = PartialStruct.Base; + } + CombinedInfo.Exprs.push_back(VD); + // Base is the base of the struct + CombinedInfo.BasePointers.push_back(PartialStruct.Base.getPointer()); + // Pointer is the address of the lowest element + llvm::Value *LB = LBAddr.getPointer(); + CombinedInfo.Pointers.push_back(LB); + // There should not be a mapper for a combined entry. + CombinedInfo.Mappers.push_back(nullptr); + // Size is (addr of {highest+1} element) - (addr of lowest element) + llvm::Value *HB = HBAddr.getPointer(); + llvm::Value *HAddr = CGF.Builder.CreateConstGEP1_32(HB, /*Idx0=*/1); + llvm::Value *CLAddr = CGF.Builder.CreatePointerCast(LB, CGF.VoidPtrTy); + llvm::Value *CHAddr = CGF.Builder.CreatePointerCast(HAddr, CGF.VoidPtrTy); + llvm::Value *Diff = CGF.Builder.CreatePtrDiff(CHAddr, CLAddr); + llvm::Value *Size = CGF.Builder.CreateIntCast(Diff, CGF.Int64Ty, + /*isSigned=*/false); + CombinedInfo.Sizes.push_back(Size); + // Map type is always TARGET_PARAM, if generate info for captures. + CombinedInfo.Types.push_back(NotTargetParams ? OMP_MAP_NONE + : OMP_MAP_TARGET_PARAM); + // If any element has the present modifier, then make sure the runtime + // doesn't attempt to allocate the struct. + if (CurTypes.end() != + llvm::find_if(CurTypes, [](OpenMPOffloadMappingFlags Type) { + return Type & OMP_MAP_PRESENT; + })) + CombinedInfo.Types.back() |= OMP_MAP_PRESENT; + // Remove TARGET_PARAM flag from the first element + (*CurTypes.begin()) &= ~OMP_MAP_TARGET_PARAM; - for (const MapInfo &L : M.second) { - assert(!L.Components.empty() && - "Not expecting declaration with no component lists."); - generateInfoForComponentList( - L.MapType, L.MapModifiers, L.MotionModifiers, L.Components, CurInfo, - PartialStruct, IsFirstComponentList, L.IsImplicit, L.Mapper, - L.ForDeviceAddr, VD, L.VarRef); - IsFirstComponentList = false; - } + // All other current entries will be MEMBER_OF the combined entry + // (except for PTR_AND_OBJ entries which do not have a placeholder value + // 0xFFFF in the MEMBER_OF field). + OpenMPOffloadMappingFlags MemberOfFlag = + getMemberOfFlag(CombinedInfo.BasePointers.size() - 1); + for (auto &M : CurTypes) + setCorrectMemberOfFlag(M, MemberOfFlag); + } - // If there is an entry in PartialStruct it means we have a struct with - // individual members mapped. Emit an extra combined entry. - if (PartialStruct.Base.isValid()) { - CurInfo.NonContigInfo.Dims.push_back(0); - emitCombinedEntry(CombinedInfo, CurInfo.Types, PartialStruct, VD); - } + /// Generate all the base pointers, section pointers, sizes, map types, and + /// mappers for the extracted mappable expressions (all included in \a + /// CombinedInfo). Also, for each item that relates with a device pointer, a + /// pair of the relevant declaration and index where it occurs is appended to + /// the device pointers info array. + void generateAllInfo( + MapCombinedInfoTy &CombinedInfo, + const llvm::DenseSet> &SkipVarSet = + llvm::DenseSet>()) const { + assert(CurDir.is() && + "Expect a executable directive"); + const auto *CurExecDir = CurDir.get(); + generateAllInfoForClauses(CurExecDir->clauses(), CombinedInfo, SkipVarSet); + } - // We need to append the results of this capture to what we already have. - CombinedInfo.append(CurInfo); - } + /// Generate all the base pointers, section pointers, sizes, map types, and + /// mappers for the extracted map clauses of user-defined mapper (all included + /// in \a CombinedInfo). + void generateAllInfoForMapper(MapCombinedInfoTy &CombinedInfo) const { + assert(CurDir.is() && + "Expect a declare mapper directive"); + const auto *CurMapperDir = CurDir.get(); + generateAllInfoForClauses(CurMapperDir->clauses(), CombinedInfo); } /// Emit capture info for lambdas for variables captured by reference. @@ -8806,6 +8827,25 @@ ++EI; } } + llvm::stable_sort(DeclComponentLists, [](const MapData &LHS, + const MapData &RHS) { + ArrayRef MapModifiers = std::get<2>(LHS); + OpenMPMapClauseKind MapType = std::get<1>(RHS); + bool HasPresent = !MapModifiers.empty() && + llvm::any_of(MapModifiers, [](OpenMPMapModifierKind K) { + return K == clang::OMPC_MAP_MODIFIER_present; + }); + bool HasAllocs = MapType == OMPC_MAP_alloc; + MapModifiers = std::get<2>(RHS); + MapType = std::get<1>(LHS); + bool HasPresentR = + !MapModifiers.empty() && + llvm::any_of(MapModifiers, [](OpenMPMapModifierKind K) { + return K == clang::OMPC_MAP_MODIFIER_present; + }); + bool HasAllocsR = MapType == OMPC_MAP_alloc; + return (HasPresent && !HasPresentR) || (HasAllocs && !HasAllocsR); + }); // Find overlapping elements (including the offset from the base element). llvm::SmallDenseMap< @@ -8841,11 +8881,23 @@ if (CI->getAssociatedDeclaration() != SI->getAssociatedDeclaration()) break; } - // Found overlapping if, at least for one component, reached the head of - // the components list. + // Found overlapping if, at least for one component, reached the head + // of the components list. if (CI == CE || SI == SE) { - assert((CI != CE || SI != SE) && - "Unexpected full match of the mapping components."); + // Ignore it if it is the same component. + if (CI == CE && SI == SE) + continue; + const auto It = (SI == SE) ? CI : SI; + // If one component is a pointer and another one is a kind of + // dereference of this pointer (array subscript, section, dereference, + // etc.), it is not an overlapping. + if (!isa(It->getAssociatedExpression()) || + std::prev(It) + ->getAssociatedExpression() + ->getType() + .getNonReferenceType() + ->isPointerType()) + continue; const MapData &BaseData = CI == CE ? L : L1; OMPClauseMappableExprCommon::MappableExprComponentListRef SubData = SI == SE ? Components : Components1; @@ -8866,7 +8918,7 @@ } } for (auto &Pair : OverlappedData) { - llvm::sort( + llvm::stable_sort( Pair.getSecond(), [&Layout]( OMPClauseMappableExprCommon::MappableExprComponentListRef First, @@ -8908,6 +8960,7 @@ // Associated with a capture, because the mapping flags depend on it. // Go through all of the elements with the overlapped elements. + bool IsFirstComponentList = true; for (const auto &Pair : OverlappedData) { const MapData &L = *Pair.getFirst(); OMPClauseMappableExprCommon::MappableExprComponentListRef Components; @@ -8920,14 +8973,13 @@ L; ArrayRef OverlappedComponents = Pair.getSecond(); - bool IsFirstComponentList = true; generateInfoForComponentList( MapType, MapModifiers, llvm::None, Components, CombinedInfo, PartialStruct, IsFirstComponentList, IsImplicit, Mapper, /*ForDeviceAddr=*/false, VD, VarRef, OverlappedComponents); + IsFirstComponentList = false; } // Go through other elements without overlapped elements. - bool IsFirstComponentList = OverlappedData.empty(); for (const MapData &L : DeclComponentLists) { OMPClauseMappableExprCommon::MappableExprComponentListRef Components; OpenMPMapClauseKind MapType; @@ -9516,13 +9568,13 @@ /// int64_t size, int64_t type, /// void *name = nullptr) { /// // Allocate space for an array section first. -/// if (size > 1 && !maptype.IsDelete) +/// if ((size > 1 || base != begin) && !maptype.IsDelete) /// __tgt_push_mapper_component(rt_mapper_handle, base, begin, /// size*sizeof(Ty), clearToFrom(type)); /// // Map members. /// for (unsigned i = 0; i < size; i++) { /// // For each component specified by this mapper: -/// for (auto c : all_components) { +/// for (auto c : begin[i]->all_components) { /// if (c.hasMapper()) /// (*c.Mapper())(rt_mapper_handle, c.arg_base, c.arg_begin, c.arg_size, /// c.arg_type, c.arg_name); @@ -9533,7 +9585,7 @@ /// } /// } /// // Delete the array section. -/// if (size > 1 && maptype.IsDelete) +/// if ((size > 1 || base != begin) && maptype.IsDelete) /// __tgt_push_mapper_component(rt_mapper_handle, base, begin, /// size*sizeof(Ty), clearToFrom(type)); /// } @@ -9585,20 +9637,10 @@ // Start the mapper function code generation. CodeGenFunction MapperCGF(CGM); MapperCGF.StartFunction(GlobalDecl(), C.VoidTy, Fn, FnInfo, Args, Loc, Loc); - // Compute the starting and end addreses of array elements. + // Compute the starting and end addresses of array elements. llvm::Value *Size = MapperCGF.EmitLoadOfScalar( MapperCGF.GetAddrOfLocalVar(&SizeArg), /*Volatile=*/false, C.getPointerType(Int64Ty), Loc); - // Convert the size in bytes into the number of array elements. - Size = MapperCGF.Builder.CreateExactUDiv( - Size, MapperCGF.Builder.getInt64(ElementSize.getQuantity())); - llvm::Value *PtrBegin = MapperCGF.Builder.CreateBitCast( - MapperCGF.GetAddrOfLocalVar(&BeginArg).getPointer(), - CGM.getTypes().ConvertTypeForMem(C.getPointerType(PtrTy))); - llvm::Value *PtrEnd = MapperCGF.Builder.CreateGEP(PtrBegin, Size); - llvm::Value *MapType = MapperCGF.EmitLoadOfScalar( - MapperCGF.GetAddrOfLocalVar(&TypeArg), /*Volatile=*/false, - C.getPointerType(Int64Ty), Loc); // Prepare common arguments for array initiation and deletion. llvm::Value *Handle = MapperCGF.EmitLoadOfScalar( MapperCGF.GetAddrOfLocalVar(&HandleArg), @@ -9609,6 +9651,15 @@ llvm::Value *BeginIn = MapperCGF.EmitLoadOfScalar( MapperCGF.GetAddrOfLocalVar(&BeginArg), /*Volatile=*/false, C.getPointerType(C.VoidPtrTy), Loc); + // Convert the size in bytes into the number of array elements. + Size = MapperCGF.Builder.CreateExactUDiv( + Size, MapperCGF.Builder.getInt64(ElementSize.getQuantity())); + llvm::Value *PtrBegin = MapperCGF.Builder.CreateBitCast( + BeginIn, CGM.getTypes().ConvertTypeForMem(PtrTy)); + llvm::Value *PtrEnd = MapperCGF.Builder.CreateGEP(PtrBegin, Size); + llvm::Value *MapType = MapperCGF.EmitLoadOfScalar( + MapperCGF.GetAddrOfLocalVar(&TypeArg), /*Volatile=*/false, + C.getPointerType(Int64Ty), Loc); // Emit array initiation if this is an array section and \p MapType indicates // that memory allocation is required. @@ -9640,11 +9691,7 @@ .alignmentOfArrayElement(ElementSize)); // Privatize the declared variable of mapper to be the current array element. CodeGenFunction::OMPPrivateScope Scope(MapperCGF); - Scope.addPrivate(MapperVarDecl, [&MapperCGF, PtrCurrent, PtrTy]() { - return MapperCGF - .EmitLoadOfPointerLValue(PtrCurrent, PtrTy->castAs()) - .getAddress(MapperCGF); - }); + Scope.addPrivate(MapperVarDecl, [PtrCurrent]() { return PtrCurrent; }); (void)Scope.Privatize(); // Get map clause information. Fill up the arrays with all mapped variables. @@ -9676,28 +9723,9 @@ : emitMappingInformation(MapperCGF, OMPBuilder, Info.Exprs[I]); // Extract the MEMBER_OF field from the map type. - llvm::BasicBlock *MemberBB = MapperCGF.createBasicBlock("omp.member"); - MapperCGF.EmitBlock(MemberBB); llvm::Value *OriMapType = MapperCGF.Builder.getInt64(Info.Types[I]); - llvm::Value *Member = MapperCGF.Builder.CreateAnd( - OriMapType, - MapperCGF.Builder.getInt64(MappableExprsHandler::OMP_MAP_MEMBER_OF)); - llvm::BasicBlock *MemberCombineBB = - MapperCGF.createBasicBlock("omp.member.combine"); - llvm::BasicBlock *TypeBB = MapperCGF.createBasicBlock("omp.type"); - llvm::Value *IsMember = MapperCGF.Builder.CreateIsNull(Member); - MapperCGF.Builder.CreateCondBr(IsMember, TypeBB, MemberCombineBB); - // Add the number of pre-existing components to the MEMBER_OF field if it - // is valid. - MapperCGF.EmitBlock(MemberCombineBB); - llvm::Value *CombinedMember = + llvm::Value *MemberMapType = MapperCGF.Builder.CreateNUWAdd(OriMapType, ShiftedPreviousSize); - // Do nothing if it is not a member of previous components. - MapperCGF.EmitBlock(TypeBB); - llvm::PHINode *MemberMapType = - MapperCGF.Builder.CreatePHI(CGM.Int64Ty, 4, "omp.membermaptype"); - MemberMapType->addIncoming(OriMapType, MemberBB); - MemberMapType->addIncoming(CombinedMember, MemberCombineBB); // Combine the map type inherited from user-defined mapper with that // specified in the program. According to the OMP_MAP_TO and OMP_MAP_FROM @@ -9819,16 +9847,14 @@ StringRef Prefix = IsInit ? ".init" : ".del"; // Evaluate if this is an array section. - llvm::BasicBlock *IsDeleteBB = - MapperCGF.createBasicBlock(getName({"omp.array", Prefix, ".evaldelete"})); llvm::BasicBlock *BodyBB = MapperCGF.createBasicBlock(getName({"omp.array", Prefix})); - llvm::Value *IsArray = MapperCGF.Builder.CreateICmpSGE( + llvm::Value *IsArray = MapperCGF.Builder.CreateICmpSGT( Size, MapperCGF.Builder.getInt64(1), "omp.arrayinit.isarray"); - MapperCGF.Builder.CreateCondBr(IsArray, IsDeleteBB, ExitBB); - - // Evaluate if we are going to delete this section. - MapperCGF.EmitBlock(IsDeleteBB); + // base != begin? + llvm::Value *BaseIsBegin = MapperCGF.Builder.CreateIsNotNull( + MapperCGF.Builder.CreatePtrDiff(Base, Begin)); + llvm::Value *Cond = MapperCGF.Builder.CreateOr(IsArray, BaseIsBegin); llvm::Value *DeleteBit = MapperCGF.Builder.CreateAnd( MapType, MapperCGF.Builder.getInt64(MappableExprsHandler::OMP_MAP_DELETE)); @@ -9840,7 +9866,8 @@ DeleteCond = MapperCGF.Builder.CreateIsNotNull( DeleteBit, getName({"omp.array", Prefix, ".delete"})); } - MapperCGF.Builder.CreateCondBr(DeleteCond, BodyBB, ExitBB); + Cond = MapperCGF.Builder.CreateAnd(Cond, DeleteCond); + MapperCGF.Builder.CreateCondBr(Cond, BodyBB, ExitBB); MapperCGF.EmitBlock(BodyBB); // Get the array size by multiplying element size and element number (i.e., \p @@ -10097,7 +10124,7 @@ llvm::DenseSet> MappedVarSet; auto RI = CS.getCapturedRecordDecl()->field_begin(); - auto CV = CapturedVars.begin(); + auto *CV = CapturedVars.begin(); for (CapturedStmt::const_capture_iterator CI = CS.capture_begin(), CE = CS.capture_end(); CI != CE; ++CI, ++RI, ++CV) { diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -3564,9 +3564,11 @@ !Stack->isLoopControlVariable(VD).first) { if (!Stack->checkMappableExprComponentListsForDecl( VD, /*CurrentRegionOnly=*/true, - [](OMPClauseMappableExprCommon::MappableExprComponentListRef - StackComponents, - OpenMPClauseKind) { + [this](OMPClauseMappableExprCommon::MappableExprComponentListRef + StackComponents, + OpenMPClauseKind) { + if (SemaRef.LangOpts.OpenMP >= 50) + return !StackComponents.empty(); // Variable is used if it has been marked as an array, array // section, array shaping or the variable iself. return StackComponents.size() == 1 || @@ -17579,7 +17581,9 @@ ERange, CKind, &EnclosingExpr, CurComponents](OMPClauseMappableExprCommon::MappableExprComponentListRef StackComponents, - OpenMPClauseKind) { + OpenMPClauseKind Kind) { + if (CKind == Kind && SemaRef.LangOpts.OpenMP >= 50) + return false; assert(!StackComponents.empty() && "Map clause expression with no components!"); assert(StackComponents.back().getAssociatedDeclaration() == VD && @@ -18133,8 +18137,6 @@ DSAS, Type)) continue; - Type = I->getAssociatedDeclaration()->getType().getNonReferenceType(); - if (CKind == OMPC_map) { // target enter data // OpenMP [2.10.2, Restrictions, p. 99] diff --git a/clang/test/OpenMP/declare_mapper_codegen.cpp b/clang/test/OpenMP/declare_mapper_codegen.cpp --- a/clang/test/OpenMP/declare_mapper_codegen.cpp +++ b/clang/test/OpenMP/declare_mapper_codegen.cpp @@ -97,17 +97,21 @@ // CK0-32-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], 8 // CK0-DAG: [[TYPE:%.+]] = load i64, i64* [[TYPEADDR]] // CK0-DAG: [[HANDLE:%.+]] = load i8*, i8** [[HANDLEADDR]] -// CK0-DAG: [[PTRBEGIN:%.+]] = bitcast i8** [[VPTRADDR]] to %class.C** -// CK0-DAG: [[PTREND:%.+]] = getelementptr %class.C*, %class.C** [[PTRBEGIN]], i64 [[SIZE]] // CK0-DAG: [[BPTR:%.+]] = load i8*, i8** [[BPTRADDR]] // CK0-DAG: [[BEGIN:%.+]] = load i8*, i8** [[VPTRADDR]] -// CK0: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1 -// CK0: br i1 [[ISARRAY]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]] - -// CK0: [[INITEVALDEL]] -// CK0: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 -// CK0: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0 -// CK0: br i1 [[ISNOTDEL]], label %[[INIT:[^,]+]], label %[[LHEAD:[^,]+]] +// CK0-DAG: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1 +// CK0-DAG: [[PTRBEGIN:%.+]] = bitcast i8* [[BEGIN]] to %class.C* +// CK0-DAG: [[PTREND:%.+]] = getelementptr %class.C, %class.C* [[PTRBEGIN]], i64 [[SIZE]] +// CK0-DAG: [[BPTRI:%.+]] = ptrtoint i8* [[BPTR]] to i64 +// CK0-DAG: [[PTRI:%.+]] = ptrtoint i8* [[BEGIN]] to i64 +// CK0-DAG: [[DIF:%.+]] = sub i64 [[BPTRI]], [[PTRI]] +// CK0-DAG: [[NORM:%.+]] = sdiv exact i64 [[DIF]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) +// CK0-DAG: [[PTRSNE:%.+]] = icmp ne i64 [[NORM]], 0 +// CK0-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[PTRSNE]] +// CK0-DAG: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 +// CK0-DAG: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0 +// CK0-DAG: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]] +// CK0: br i1 [[CMP1]], label %[[INIT:[^,]+]], label %[[LHEAD:[^,]+]] // CK0: [[INIT]] // CK0-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16 // CK0-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8 @@ -116,14 +120,13 @@ // CK0: br label %[[LHEAD:[^,]+]] // CK0: [[LHEAD]] -// CK0: [[ISEMPTY:%.+]] = icmp eq %class.C** [[PTRBEGIN]], [[PTREND]] +// CK0: [[ISEMPTY:%.+]] = icmp eq %class.C* [[PTRBEGIN]], [[PTREND]] // CK0: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]] // CK0: [[LBODY]] -// CK0: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %[[LHEAD]] ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ] -// CK0: [[OBJ:%.+]] = load %class.C*, %class.C** [[PTR]] -// CK0-DAG: [[ABEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 0 -// CK0-DAG: [[BBEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1 -// CK0-DAG: [[BBEGIN2:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1 +// CK0: [[PTR:%.+]] = phi %class.C* [ [[PTRBEGIN]], %{{.+}} ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ] +// CK0-DAG: [[ABEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[PTR]], i32 0, i32 0 +// CK0-DAG: [[BBEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[PTR]], i32 0, i32 1 +// CK0-DAG: [[BBEGIN2:%.+]] = getelementptr inbounds %class.C, %class.C* [[PTR]], i32 0, i32 1 // CK0-DAG: [[BARRBEGIN:%.+]] = load double*, double** [[BBEGIN2]] // CK0-DAG: [[BARRBEGINGEP:%.+]] = getelementptr inbounds double, double* [[BARRBEGIN]], i[[sz:64|32]] 0 // CK0-DAG: [[BEND:%.+]] = getelementptr double*, double** [[BBEGIN]], i32 1 @@ -133,18 +136,11 @@ // CK0-DAG: [[BENDI:%.+]] = ptrtoint i8* [[BENDV]] to i64 // CK0-DAG: [[CSIZE:%.+]] = sub i64 [[BENDI]], [[ABEGINI]] // CK0-DAG: [[CUSIZE:%.+]] = sdiv exact i64 [[CSIZE]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) -// CK0-DAG: [[BPTRADDR0BC:%.+]] = bitcast %class.C* [[OBJ]] to i8* -// CK0-DAG: [[PTRADDR0BC:%.+]] = bitcast i32* [[ABEGIN]] to i8* // CK0-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(i8* [[HANDLE]]) // CK0-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 48 -// CK0-DAG: br label %[[MEMBER:[^,]+]] -// CK0-DAG: [[MEMBER]] -// CK0-DAG: br i1 true, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]] -// CK0-DAG: [[MEMBERCOM]] -// CK0-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 0, [[SHIPRESIZE]] -// CK0-DAG: br label %[[LTYPE]] -// CK0-DAG: [[LTYPE]] -// CK0-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 0, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ] +// CK0-DAG: [[BPTRADDR0BC:%.+]] = bitcast %class.C* [[PTR]] to i8* +// CK0-DAG: [[PTRADDR0BC:%.+]] = bitcast i32* [[ABEGIN]] to i8* +// CK0-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 0, [[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:[^,]+]] @@ -166,17 +162,10 @@ // CK0-DAG: [[TYEND]] // CK0-DAG: [[PHITYPE0:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ] // CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR0BC]], i8* [[PTRADDR0BC]], i64 [[CUSIZE]], i64 [[PHITYPE0]], {{.*}}) -// CK0-DAG: [[BPTRADDR1BC:%.+]] = bitcast %class.C* [[OBJ]] to i8* +// CK0-DAG: [[BPTRADDR1BC:%.+]] = bitcast %class.C* [[PTR]] to i8* // CK0-DAG: [[PTRADDR1BC:%.+]] = bitcast i32* [[ABEGIN]] to i8* -// CK0-DAG: br label %[[MEMBER:[^,]+]] -// CK0-DAG: [[MEMBER]] -// CK0-DAG: br i1 false, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]] -// CK0-DAG: [[MEMBERCOM]] // 281474976710659 == 0x1,000,000,003 -// CK0-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 281474976710659, [[SHIPRESIZE]] -// CK0-DAG: br label %[[LTYPE]] -// CK0-DAG: [[LTYPE]] -// CK0-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 281474976710659, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ] +// CK0-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710659, [[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:[^,]+]] @@ -200,15 +189,8 @@ // CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR1BC]], i8* [[PTRADDR1BC]], i64 4, i64 [[TYPE1]], {{.*}}) // CK0-DAG: [[BPTRADDR2BC:%.+]] = bitcast double** [[BBEGIN]] to i8* // CK0-DAG: [[PTRADDR2BC:%.+]] = bitcast double* [[BARRBEGINGEP]] to i8* -// CK0-DAG: br label %[[MEMBER:[^,]+]] -// CK0-DAG: [[MEMBER]] -// CK0-DAG: br i1 false, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]] -// CK0-DAG: [[MEMBERCOM]] // 281474976710675 == 0x1,000,000,013 -// CK0-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 281474976710675, [[SHIPRESIZE]] -// CK0-DAG: br label %[[LTYPE]] -// CK0-DAG: [[LTYPE]] -// CK0-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 281474976710675, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ] +// CK0-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710675, [[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:[^,]+]] @@ -230,18 +212,23 @@ // CK0-DAG: [[TYEND]] // CK0-DAG: [[TYPE2:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ] // CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR2BC]], i8* [[PTRADDR2BC]], i64 16, i64 [[TYPE2]], {{.*}}) -// CK0: [[PTRNEXT]] = getelementptr %class.C*, %class.C** [[PTR]], i32 1 -// CK0: [[ISDONE:%.+]] = icmp eq %class.C** [[PTRNEXT]], [[PTREND]] +// CK0: [[PTRNEXT]] = getelementptr %class.C, %class.C* [[PTR]], i32 1 +// CK0: [[ISDONE:%.+]] = icmp eq %class.C* [[PTRNEXT]], [[PTREND]] // CK0: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]] // CK0: [[LEXIT]] -// CK0: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1 -// CK0: br i1 [[ISARRAY]], label %[[EVALDEL:[^,]+]], label %[[DONE]] -// CK0: [[EVALDEL]] +// CK0: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1 +// CK0: [[BPTRI:%.+]] = ptrtoint i8* [[BPTR]] to i64 +// CK0: [[PTRI:%.+]] = ptrtoint i8* [[BEGIN]] to i64 +// CK0: [[DIF:%.+]] = sub i64 [[BPTRI]], [[PTRI]] +// CK0: [[NORM:%.+]] = sdiv exact i64 [[DIF]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) +// CK0: [[PTRSNE:%.+]] = icmp ne i64 [[NORM]], 0 +// CK0: [[CMP:%.+]] = or i1 [[ISARRAY]], [[PTRSNE]] // CK0: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 -// CK0: [[ISDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0 -// CK0: br i1 [[ISDEL]], label %[[DEL:[^,]+]], label %[[DONE]] -// CK0: [[DEL]] +// CK0: [[ISNOTDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0 +// CK0: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]] +// CK0: br i1 [[CMP1]], label %[[EVALDEL:[^,]+]], label %[[DONE]] +// CK0: [[EVALDEL]] // CK0-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16 // CK0-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8 // CK0-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4 @@ -662,42 +649,39 @@ // CK1-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], 4 // CK1-DAG: [[TYPE:%.+]] = load i64, i64* [[TYPEADDR]] // CK1-DAG: [[HANDLE:%.+]] = load i8*, i8** [[HANDLEADDR]] -// CK1-DAG: [[PTRBEGIN:%.+]] = bitcast i8** [[VPTRADDR]] to %class.C** -// CK1-DAG: [[PTREND:%.+]] = getelementptr %class.C*, %class.C** [[PTRBEGIN]], i64 [[SIZE]] // CK1-DAG: [[BPTR:%.+]] = load i8*, i8** [[BPTRADDR]] // CK1-DAG: [[BEGIN:%.+]] = load i8*, i8** [[VPTRADDR]] -// CK1: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1 -// CK1: br i1 [[ISARRAY]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]] +// CK1-DAG: [[PTRBEGIN:%.+]] = bitcast i8* [[BEGIN]] to %class.C* +// CK1-DAG: [[PTREND:%.+]] = getelementptr %class.C, %class.C* [[PTRBEGIN]], i64 [[SIZE]] +// CK1-DAG: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1 +// CK1-DAG: [[BPTRI:%.+]] = ptrtoint i8* [[BPTR]] to i64 +// CK1-DAG: [[PTRI:%.+]] = ptrtoint i8* [[BEGIN]] to i64 +// CK1-DAG: [[DIF:%.+]] = sub i64 [[BPTRI]], [[PTRI]] +// CK1-DAG: [[NORM:%.+]] = sdiv exact i64 [[DIF]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) +// CK1-DAG: [[PTRSNE:%.+]] = icmp ne i64 [[NORM]], 0 +// CK1-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[PTRSNE]] +// CK1-DAG: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 +// CK1-DAG: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0 +// CK1-DAG: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]] +// CK1: br i1 [[CMP1]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]] // CK1: [[INITEVALDEL]] -// CK1: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 -// CK1: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0 -// CK1: br i1 [[ISNOTDEL]], label %[[INIT:[^,]+]], label %[[LHEAD:[^,]+]] -// CK1: [[INIT]] // CK1-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 4 // CK1-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4 // CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]], {{.*}}) // CK1: br label %[[LHEAD:[^,]+]] // CK1: [[LHEAD]] -// CK1: [[ISEMPTY:%.+]] = icmp eq %class.C** [[PTRBEGIN]], [[PTREND]] +// CK1: [[ISEMPTY:%.+]] = icmp eq %class.C* [[PTRBEGIN]], [[PTREND]] // CK1: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]] // CK1: [[LBODY]] -// CK1: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %[[LHEAD]] ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ] -// CK1: [[OBJ:%.+]] = load %class.C*, %class.C** [[PTR]] -// CK1-DAG: [[ABEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 0 +// CK1: [[PTR:%.+]] = phi %class.C* [ [[PTRBEGIN]], %{{.+}} ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ] +// CK1-DAG: [[ABEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[PTR]], i32 0, i32 0 // CK1-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(i8* [[HANDLE]]) // CK1-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 48 -// CK1-DAG: [[BPTRADDR1BC:%.+]] = bitcast %class.C* [[OBJ]] to i8* +// CK1-DAG: [[BPTRADDR1BC:%.+]] = bitcast %class.C* [[PTR]] to i8* // CK1-DAG: [[PTRADDR1BC:%.+]] = bitcast i32* [[ABEGIN]] to i8* -// CK1-DAG: br label %[[MEMBER:[^,]+]] -// CK1-DAG: [[MEMBER]] -// CK1-DAG: br i1 true, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]] -// CK1-DAG: [[MEMBERCOM]] -// CK1-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 35, [[SHIPRESIZE]] -// CK1-DAG: br label %[[LTYPE]] -// CK1-DAG: [[LTYPE]] -// CK1-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 35, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ] +// CK1-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 3, [[SHIPRESIZE]] // CK1-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3 // CK1-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0 // CK1-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]] @@ -719,18 +703,21 @@ // CK1-DAG: [[TYEND]] // CK1-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ] // CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR1BC]], i8* [[PTRADDR1BC]], i64 4, i64 [[TYPE1]], {{.*}}) -// CK1: [[PTRNEXT]] = getelementptr %class.C*, %class.C** [[PTR]], i32 1 -// CK1: [[ISDONE:%.+]] = icmp eq %class.C** [[PTRNEXT]], [[PTREND]] +// CK1: [[PTRNEXT]] = getelementptr %class.C, %class.C* [[PTR]], i32 1 +// CK1: [[ISDONE:%.+]] = icmp eq %class.C* [[PTRNEXT]], [[PTREND]] // CK1: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]] // CK1: [[LEXIT]] -// CK1: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1 -// CK1: br i1 [[ISARRAY]], label %[[EVALDEL:[^,]+]], label %[[DONE]] -// CK1: [[EVALDEL]] +// CK1: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1 +// CK1: [[BPTRI:%.+]] = ptrtoint i8* [[BPTR]] to i64 +// CK1: [[PTRI:%.+]] = ptrtoint i8* [[BEGIN]] to i64 +// CK1: [[DIF:%.+]] = sub i64 [[BPTRI]], [[PTRI]] +// CK1: [[NORM:%.+]] = sdiv exact i64 [[DIF]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) +// CK1: [[PTRSNE:%.+]] = icmp ne i64 [[NORM]], 0 +// CK1: [[CMP:%.+]] = or i1 [[ISARRAY]], [[PTRSNE]] // CK1: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 -// CK1: [[ISDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0 -// CK1: br i1 [[ISDEL]], label %[[DEL:[^,]+]], label %[[DONE]] -// CK1: [[DEL]] +// CK1: [[ISNOTDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0 +// CK1: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]] // CK1-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 4 // CK1-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4 // CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]], {{.*}}) @@ -786,42 +773,39 @@ // CK2-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], 16 // CK2-DAG: [[TYPE:%.+]] = load i64, i64* [[TYPEADDR]] // CK2-DAG: [[HANDLE:%.+]] = load i8*, i8** [[HANDLEADDR]] -// CK2-DAG: [[PTRBEGIN:%.+]] = bitcast i8** [[VPTRADDR]] to %class.C** -// CK2-DAG: [[PTREND:%.+]] = getelementptr %class.C*, %class.C** [[PTRBEGIN]], i64 [[SIZE]] // CK2-DAG: [[BPTR:%.+]] = load i8*, i8** [[BPTRADDR]] // CK2-DAG: [[BEGIN:%.+]] = load i8*, i8** [[VPTRADDR]] -// CK2: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1 -// CK2: br i1 [[ISARRAY]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]] +// CK2-DAG: [[PTRBEGIN:%.+]] = bitcast i8* [[BEGIN]] to %class.C* +// CK2-DAG: [[PTREND:%.+]] = getelementptr %class.C, %class.C* [[PTRBEGIN]], i64 [[SIZE]] +// CK2-DAG: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1 +// CK2-DAG: [[BPTRI:%.+]] = ptrtoint i8* [[BPTR]] to i64 +// CK2-DAG: [[PTRI:%.+]] = ptrtoint i8* [[BEGIN]] to i64 +// CK2-DAG: [[DIF:%.+]] = sub i64 [[BPTRI]], [[PTRI]] +// CK2-DAG: [[NORM:%.+]] = sdiv exact i64 [[DIF]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) +// CK2-DAG: [[PTRSNE:%.+]] = icmp ne i64 [[NORM]], 0 +// CK2-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[PTRSNE]] +// CK2-DAG: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 +// CK2-DAG: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0 +// CK2-DAG: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]] +// CK2: br i1 [[CMP1]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]] // CK2: [[INITEVALDEL]] -// CK2: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 -// CK2: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0 -// CK2: br i1 [[ISNOTDEL]], label %[[INIT:[^,]+]], label %[[LHEAD:[^,]+]] -// CK2: [[INIT]] // CK2-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16 // CK2-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4 // CK2: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]], {{.*}}) // CK2: br label %[[LHEAD:[^,]+]] // CK2: [[LHEAD]] -// CK2: [[ISEMPTY:%.+]] = icmp eq %class.C** [[PTRBEGIN]], [[PTREND]] +// CK2: [[ISEMPTY:%.+]] = icmp eq %class.C* [[PTRBEGIN]], [[PTREND]] // CK2: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]] // CK2: [[LBODY]] -// CK2: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %[[LHEAD]] ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ] -// CK2: [[OBJ:%.+]] = load %class.C*, %class.C** [[PTR]] -// CK2-DAG: [[BBEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1 +// CK2: [[PTR:%.+]] = phi %class.C* [ [[PTRBEGIN]], %{{.+}} ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ] +// CK2-DAG: [[BBEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[PTR]], i32 0, i32 1 // CK2-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(i8* [[HANDLE]]) // CK2-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 48 -// CK2-DAG: [[BPTRADDR1BC:%.+]] = bitcast %class.C* [[OBJ]] to i8* +// CK2-DAG: [[BPTRADDR1BC:%.+]] = bitcast %class.C* [[PTR]] to i8* // CK2-DAG: [[PTRADDR1BC:%.+]] = bitcast %class.B* [[BBEGIN]] to i8* -// CK2-DAG: br label %[[MEMBER:[^,]+]] -// CK2-DAG: [[MEMBER]] -// CK2-DAG: br i1 true, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]] -// CK2-DAG: [[MEMBERCOM]] -// CK2-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 35, [[SHIPRESIZE]] -// CK2-DAG: br label %[[LTYPE]] -// CK2-DAG: [[LTYPE]] -// CK2-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 35, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ] +// CK2-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 3, [[SHIPRESIZE]] // CK2-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3 // CK2-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0 // CK2-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]] @@ -843,18 +827,23 @@ // CK2-DAG: [[TYEND]] // CK2-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ] // CK2: call void [[BMPRFUNC]](i8* [[HANDLE]], i8* [[BPTRADDR1BC]], i8* [[PTRADDR1BC]], i64 8, i64 [[TYPE1]], {{.*}}) -// CK2: [[PTRNEXT]] = getelementptr %class.C*, %class.C** [[PTR]], i32 1 -// CK2: [[ISDONE:%.+]] = icmp eq %class.C** [[PTRNEXT]], [[PTREND]] +// CK2: [[PTRNEXT]] = getelementptr %class.C, %class.C* [[PTR]], i32 1 +// CK2: [[ISDONE:%.+]] = icmp eq %class.C* [[PTRNEXT]], [[PTREND]] // CK2: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]] // CK2: [[LEXIT]] -// CK2: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1 -// CK2: br i1 [[ISARRAY]], label %[[EVALDEL:[^,]+]], label %[[DONE]] -// CK2: [[EVALDEL]] +// CK2: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1 +// CK2: [[BPTRI:%.+]] = ptrtoint i8* [[BPTR]] to i64 +// CK2: [[PTRI:%.+]] = ptrtoint i8* [[BEGIN]] to i64 +// CK2: [[DIF:%.+]] = sub i64 [[BPTRI]], [[PTRI]] +// CK2: [[NORM:%.+]] = sdiv exact i64 [[DIF]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) +// CK2: [[PTRSNE:%.+]] = icmp ne i64 [[NORM]], 0 +// CK2: [[CMP:%.+]] = or i1 [[ISARRAY]], [[PTRSNE]] // CK2: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 -// CK2: [[ISDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0 -// CK2: br i1 [[ISDEL]], label %[[DEL:[^,]+]], label %[[DONE]] -// CK2: [[DEL]] +// CK2: [[ISNOTDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0 +// CK2: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]] +// CK2: br i1 [[CMP1]], label %[[EVALDEL:[^,]+]], label %[[DONE]] +// CK2: [[EVALDEL]] // CK2-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16 // CK2-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4 // CK2: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]], {{.*}}) @@ -991,18 +980,23 @@ // CK4-32-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], 8 // CK4-DAG: [[TYPE:%.+]] = load i64, i64* [[TYPEADDR]] // CK4-DAG: [[HANDLE:%.+]] = load i8*, i8** [[HANDLEADDR]] -// CK4-DAG: [[PTRBEGIN:%.+]] = bitcast i8** [[VPTRADDR]] to %class.C** -// CK4-DAG: [[PTREND:%.+]] = getelementptr %class.C*, %class.C** [[PTRBEGIN]], i64 [[SIZE]] // CK4-DAG: [[BPTR:%.+]] = load i8*, i8** [[BPTRADDR]] // CK4-DAG: [[BEGIN:%.+]] = load i8*, i8** [[VPTRADDR]] -// CK4: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1 -// CK4: br i1 [[ISARRAY]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]] +// CK4-DAG: [[PTRBEGIN:%.+]] = bitcast i8* [[BEGIN]] to %class.C* +// CK4-DAG: [[PTREND:%.+]] = getelementptr %class.C, %class.C* [[PTRBEGIN]], i64 [[SIZE]] +// CK4-DAG: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1 +// CK4-DAG: [[BPTRI:%.+]] = ptrtoint i8* [[BPTR]] to i64 +// CK4-DAG: [[PTRI:%.+]] = ptrtoint i8* [[BEGIN]] to i64 +// CK4-DAG: [[DIF:%.+]] = sub i64 [[BPTRI]], [[PTRI]] +// CK4-DAG: [[NORM:%.+]] = sdiv exact i64 [[DIF]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) +// CK4-DAG: [[PTRSNE:%.+]] = icmp ne i64 [[NORM]], 0 +// CK4-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[PTRSNE]] +// CK4-DAG: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 +// CK4-DAG: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0 +// CK4-DAG: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]] +// CK4: br i1 [[CMP1]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]] // CK4: [[INITEVALDEL]] -// CK4: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 -// CK4: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0 -// CK4: br i1 [[ISNOTDEL]], label %[[INIT:[^,]+]], label %[[LHEAD:[^,]+]] -// CK4: [[INIT]] // CK4-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16 // CK4-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8 // CK4-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4 @@ -1010,14 +1004,13 @@ // CK4: br label %[[LHEAD:[^,]+]] // CK4: [[LHEAD]] -// CK4: [[ISEMPTY:%.+]] = icmp eq %class.C** [[PTRBEGIN]], [[PTREND]] +// CK4: [[ISEMPTY:%.+]] = icmp eq %class.C* [[PTRBEGIN]], [[PTREND]] // CK4: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]] // CK4: [[LBODY]] -// CK4: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %[[LHEAD]] ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ] -// CK4: [[OBJ:%.+]] = load %class.C*, %class.C** [[PTR]] -// CK4-DAG: [[ABEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 0 -// CK4-DAG: [[BBEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1 -// CK4-DAG: [[BBEGIN2:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1 +// CK4: [[PTR:%.+]] = phi %class.C* [ [[PTRBEGIN]], %{{.+}} ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ] +// CK4-DAG: [[ABEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[PTR]], i32 0, i32 0 +// CK4-DAG: [[BBEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[PTR]], i32 0, i32 1 +// CK4-DAG: [[BBEGIN2:%.+]] = getelementptr inbounds %class.C, %class.C* [[PTR]], i32 0, i32 1 // CK4-DAG: [[BARRBEGIN:%.+]] = load double*, double** [[BBEGIN2]] // CK4-DAG: [[BARRBEGINGEP:%.+]] = getelementptr inbounds double, double* [[BARRBEGIN]], i[[sz:64|32]] 0 // CK4-DAG: [[BEND:%.+]] = getelementptr double*, double** [[BBEGIN]], i32 1 @@ -1027,18 +1020,11 @@ // CK4-DAG: [[BENDI:%.+]] = ptrtoint i8* [[BENDV]] to i64 // CK4-DAG: [[CSIZE:%.+]] = sub i64 [[BENDI]], [[ABEGINI]] // CK4-DAG: [[CUSIZE:%.+]] = sdiv exact i64 [[CSIZE]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) -// CK4-DAG: [[BPTRADDR0BC:%.+]] = bitcast %class.C* [[OBJ]] to i8* +// CK4-DAG: [[BPTRADDR0BC:%.+]] = bitcast %class.C* [[PTR]] to i8* // CK4-DAG: [[PTRADDR0BC:%.+]] = bitcast i32* [[ABEGIN]] to i8* // CK4-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(i8* [[HANDLE]]) // CK4-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 48 -// CK4-DAG: br label %[[MEMBER:[^,]+]] -// CK4-DAG: [[MEMBER]] -// CK4-DAG: br i1 true, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]] -// CK4-DAG: [[MEMBERCOM]] -// CK4-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 0, [[SHIPRESIZE]] -// CK4-DAG: br label %[[LTYPE]] -// CK4-DAG: [[LTYPE]] -// CK4-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 0, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ] +// CK4-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 0, [[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:[^,]+]] @@ -1060,17 +1046,10 @@ // CK4-DAG: [[TYEND]] // CK4-DAG: [[PHITYPE0:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ] // CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR0BC]], i8* [[PTRADDR0BC]], i64 [[CUSIZE]], i64 [[PHITYPE0]], {{.*}}) -// CK4-DAG: [[BPTRADDR1BC:%.+]] = bitcast %class.C* [[OBJ]] to i8* +// CK4-DAG: [[BPTRADDR1BC:%.+]] = bitcast %class.C* [[PTR]] to i8* // CK4-DAG: [[PTRADDR1BC:%.+]] = bitcast i32* [[ABEGIN]] to i8* -// CK4-DAG: br label %[[MEMBER:[^,]+]] -// CK4-DAG: [[MEMBER]] -// CK4-DAG: br i1 false, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]] -// CK4-DAG: [[MEMBERCOM]] // 281474976710659 == 0x1,000,000,003 -// CK4-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 281474976710659, [[SHIPRESIZE]] -// CK4-DAG: br label %[[LTYPE]] -// CK4-DAG: [[LTYPE]] -// CK4-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 281474976710659, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ] +// CK4-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710659, [[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:[^,]+]] @@ -1094,15 +1073,8 @@ // CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR1BC]], i8* [[PTRADDR1BC]], i64 4, i64 [[TYPE1]], {{.*}}) // CK4-DAG: [[BPTRADDR2BC:%.+]] = bitcast double** [[BBEGIN]] to i8* // CK4-DAG: [[PTRADDR2BC:%.+]] = bitcast double* [[BARRBEGINGEP]] to i8* -// CK4-DAG: br label %[[MEMBER:[^,]+]] -// CK4-DAG: [[MEMBER]] -// CK4-DAG: br i1 false, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]] -// CK4-DAG: [[MEMBERCOM]] // 281474976710675 == 0x1,000,000,013 -// CK4-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 281474976710675, [[SHIPRESIZE]] -// CK4-DAG: br label %[[LTYPE]] -// CK4-DAG: [[LTYPE]] -// CK4-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 281474976710675, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ] +// CK4-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710675, [[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:[^,]+]] @@ -1124,18 +1096,23 @@ // CK4-DAG: [[TYEND]] // CK4-DAG: [[TYPE2:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ] // CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR2BC]], i8* [[PTRADDR2BC]], i64 16, i64 [[TYPE2]], {{.*}}) -// CK4: [[PTRNEXT]] = getelementptr %class.C*, %class.C** [[PTR]], i32 1 -// CK4: [[ISDONE:%.+]] = icmp eq %class.C** [[PTRNEXT]], [[PTREND]] +// CK4: [[PTRNEXT]] = getelementptr %class.C, %class.C* [[PTR]], i32 1 +// CK4: [[ISDONE:%.+]] = icmp eq %class.C* [[PTRNEXT]], [[PTREND]] // CK4: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]] // CK4: [[LEXIT]] -// CK4: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1 -// CK4: br i1 [[ISARRAY]], label %[[EVALDEL:[^,]+]], label %[[DONE]] -// CK4: [[EVALDEL]] +// CK4: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1 +// CK4: [[BPTRI:%.+]] = ptrtoint i8* [[BPTR]] to i64 +// CK4: [[PTRI:%.+]] = ptrtoint i8* [[BEGIN]] to i64 +// CK4: [[DIF:%.+]] = sub i64 [[BPTRI]], [[PTRI]] +// CK4: [[NORM:%.+]] = sdiv exact i64 [[DIF]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) +// CK4: [[PTRSNE:%.+]] = icmp ne i64 [[NORM]], 0 +// CK4: [[CMP:%.+]] = or i1 [[ISARRAY]], [[PTRSNE]] // CK4: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 -// CK4: [[ISDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0 -// CK4: br i1 [[ISDEL]], label %[[DEL:[^,]+]], label %[[DONE]] -// CK4: [[DEL]] +// CK4: [[ISNOTDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0 +// CK4: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]] +// CK4: br i1 [[CMP1]], label %[[EVALDEL:[^,]+]], label %[[DONE]] +// CK4: [[EVALDEL]] // CK4-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16 // CK4-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8 // CK4-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4 diff --git a/clang/test/OpenMP/target_data_codegen.cpp b/clang/test/OpenMP/target_data_codegen.cpp --- a/clang/test/OpenMP/target_data_codegen.cpp +++ b/clang/test/OpenMP/target_data_codegen.cpp @@ -561,8 +561,9 @@ void test_close_modifier(int arg) { S2 *ps; - // CK5: private unnamed_addr constant [5 x i64] [i64 1027, i64 0, i64 562949953421328, i64 16, i64 1043] - #pragma omp target data map(close,tofrom: arg, ps->ps->ps->ps->s) +// CK5: private unnamed_addr constant [5 x i64] [i64 1027, i64 0, i64 562949953421328, i64 16, i64 1043] +#pragma omp target data map(close, tofrom \ + : arg, ps->ps->ps->ps->s) { ++(arg); } @@ -585,8 +586,9 @@ // SIMD-ONLY2-NOT: {{__kmpc|__tgt}} #ifdef CK6 void test_close_modifier(int arg) { - // CK6: private unnamed_addr constant [1 x i64] [i64 1027] - #pragma omp target data map(close,tofrom: arg) +// CK6: private unnamed_addr constant [1 x i64] [i64 1027] +#pragma omp target data map(close, tofrom \ + : arg) {++arg;} } #endif @@ -642,36 +644,39 @@ // CK8: private unnamed_addr constant [11 x i64] - // ps1 - // - // PRESENT=0x1000 = 0x1000 - // MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003 - // MEMBER_OF_1=0x1000000000000 | PRESENT=0x1000 | PTR_AND_OBJ=0x10 = 0x1000000001010 - // PRESENT=0x1000 | PTR_AND_OBJ=0x10 = 0x1010 - // PRESENT=0x1000 | PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x1013 - // - // CK8-SAME: {{^}} [i64 [[#0x1000]], i64 [[#0x1000000000003]], - // CK8-SAME: {{^}} i64 [[#0x1000000001010]], i64 [[#0x1010]], i64 [[#0x1013]], - - // arg - // - // PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x1003 - // - // CK8-SAME: {{^}} i64 [[#0x1003]], - - // ps2 - // - // PRESENT=0x1000 = 0x1000 - // MEMBER_OF_7=0x7000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x7000000001003 - // MEMBER_OF_7=0x7000000000000 | PTR_AND_OBJ=0x10 = 0x7000000000010 - // PTR_AND_OBJ=0x10 = 0x10 - // PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x13 - // - // CK8-SAME: {{^}} i64 [[#0x1000]], i64 [[#0x7000000001003]], - // CK8-SAME: {{^}} i64 [[#0x7000000000010]], i64 [[#0x10]], i64 [[#0x13]]] - #pragma omp target data map(tofrom: ps1->s) \ - map(present,tofrom: arg, ps1->ps->ps->ps->s, ps2->s) \ - map(tofrom: ps2->ps->ps->ps->s) +// ps1 +// +// PRESENT=0x1000 = 0x1000 +// MEMBER_OF_1=0x1000000000000 | PRESENT=0x1000 | PTR_AND_OBJ=0x10 = 0x1000000001010 +// PRESENT=0x1000 | PTR_AND_OBJ=0x10 = 0x1010 +// PRESENT=0x1000 | PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x1013 +// MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003 +// +// CK8-SAME: {{^}} [i64 [[#0x1000]], i64 [[#0x1000000001010]], +// CK8-SAME: {{^}} i64 [[#0x1010]], i64 [[#0x1013]], i64 [[#0x1000000000003]], + +// arg +// +// PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x1003 +// +// CK8-SAME: {{^}} i64 [[#0x1003]], + +// ps2 +// +// PRESENT=0x1000 = 0x1000 +// MEMBER_OF_7=0x7000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x7000000001003 +// MEMBER_OF_7=0x7000000000000 | PTR_AND_OBJ=0x10 = 0x7000000000010 +// PTR_AND_OBJ=0x10 = 0x10 +// PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x13 +// +// CK8-SAME: {{^}} i64 [[#0x1000]], i64 [[#0x7000000001003]], +// CK8-SAME: {{^}} i64 [[#0x7000000000010]], i64 [[#0x10]], i64 [[#0x13]]] +#pragma omp target data map(tofrom \ + : ps1->s) \ + map(present, tofrom \ + : arg, ps1->ps->ps->ps->s, ps2->s) \ + map(tofrom \ + : ps2->ps->ps->ps->s) { ++(arg); } @@ -694,9 +699,10 @@ // SIMD-ONLY2-NOT: {{__kmpc|__tgt}} #ifdef CK9 void test_present_modifier(int arg) { - // PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x1003 - // CK9: private unnamed_addr constant [1 x i64] [i64 [[#0x1003]]] - #pragma omp target data map(present,tofrom: arg) +// PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x1003 +// CK9: private unnamed_addr constant [1 x i64] [i64 [[#0x1003]]] +#pragma omp target data map(present, tofrom \ + : arg) {++arg;} } #endif diff --git a/clang/test/OpenMP/target_enter_data_codegen.cpp b/clang/test/OpenMP/target_enter_data_codegen.cpp --- a/clang/test/OpenMP/target_enter_data_codegen.cpp +++ b/clang/test/OpenMP/target_enter_data_codegen.cpp @@ -267,7 +267,7 @@ // PRESENT=0x1000 | TO=0x1 = 0x1001 // CK1A: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1001]]] -// PRESENT=0x1000 | CLOSE=0x400 | ALWAYS=0x4 | TO=0x1 = 0x1425 +// PRESENT=0x1000 | CLOSE=0x400 | ALWAYS=0x4 | TO=0x1 = 0x1405 // CK1A: [[MTYPE01:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1405]]] // CK1A-LABEL: _Z3fooi diff --git a/clang/test/OpenMP/target_is_device_ptr_messages.cpp b/clang/test/OpenMP/target_is_device_ptr_messages.cpp --- a/clang/test/OpenMP/target_is_device_ptr_messages.cpp +++ b/clang/test/OpenMP/target_is_device_ptr_messages.cpp @@ -1,6 +1,8 @@ -// RUN: %clang_cc1 -std=c++11 -verify -fopenmp -ferror-limit 200 %s -Wuninitialized +// RUN: %clang_cc1 -std=c++11 -verify=expected,omp45 -fopenmp -fopenmp-version=45 -ferror-limit 200 %s -Wuninitialized +// RUN: %clang_cc1 -std=c++11 -verify=expected,omp50 -fopenmp -ferror-limit 200 %s -Wuninitialized -// RUN: %clang_cc1 -std=c++11 -verify -fopenmp-simd -ferror-limit 200 %s -Wuninitialized +// RUN: %clang_cc1 -std=c++11 -verify=expected,omp45 -fopenmp-simd -fopenmp-version=45 -ferror-limit 200 %s -Wuninitialized +// RUN: %clang_cc1 -std=c++11 -verify=expected,omp50 -fopenmp-simd -ferror-limit 200 %s -Wuninitialized struct ST { int *a; }; diff --git a/clang/test/OpenMP/target_map_codegen_29.cpp b/clang/test/OpenMP/target_map_codegen_29.cpp --- a/clang/test/OpenMP/target_map_codegen_29.cpp +++ b/clang/test/OpenMP/target_map_codegen_29.cpp @@ -73,9 +73,8 @@ // CK30-DAG: [[S_BEGIN_BC]] = ptrtoint i8* [[S_BEGIN:%.+]] to i64 // CK30-DAG: [[S_END_BC]] = ptrtoint i8* [[S_END:%.+]] to i64 // CK30-DAG: [[S_BEGIN]] = bitcast [[STRUCT]]* [[S]] to i8* -// CK30-DAG: [[S_END]] = getelementptr i8, i8* [[S_LAST:%.+]], i32 1 -// CK30-DAG: [[S_LAST]] = getelementptr i8, i8* [[S_BC:%.+]], i{{64|32}} {{55|27}} -// CK30-DAG: [[S_BC]] = bitcast [[STRUCT]]* [[S]] to i8* +// CK30-DAG: [[S_END]] = bitcast [[STRUCT]]* [[REAL_S_END:%.+]] to i8* +// CK30-DAG: [[REAL_S_END]] = getelementptr [[STRUCT]], [[STRUCT]]* [[S]], i32 1 // CK30-DAG: [[BASE_PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BASES]], i32 0, i32 1 // CK30-DAG: [[BC:%.+]] = bitcast i8** [[BASE_PTR]] to [[STRUCT]]** @@ -125,7 +124,9 @@ // CK30-DAG: [[S_PTR1_BC]] = ptrtoint i8* [[S_PTR1:%.+]] to i64 // CK30-DAG: [[S_END_BC]] = ptrtoint i8* [[S_END:%.+]] to i64 // CK30-DAG: [[S_PTR1]] = bitcast i32** [[PTR2]] to i8* -// CK30-DAG: [[S_END]] = getelementptr i8, i8* [[S_LAST]], i{{64|32}} 1 +// CK30-DAG: [[S_END]] = getelementptr i8, i8* [[S_LAST:%.+]], i{{64|32}} 1 +// CK30-DAG: [[S_LAST]] = getelementptr i8, i8* [[S_BC:%.+]], i{{64|32}} {{55|27}} +// CK30-DAG: [[S_BC]] = bitcast [[STRUCT]]* [[S]] to i8* // CK30-DAG: [[BASE_PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BASES]], i32 0, i32 4 // CK30-DAG: [[BC:%.+]] = bitcast i8** [[BASE_PTR]] to i32*** diff --git a/clang/test/OpenMP/target_map_codegen_31.cpp b/clang/test/OpenMP/target_map_codegen_31.cpp --- a/clang/test/OpenMP/target_map_codegen_31.cpp +++ b/clang/test/OpenMP/target_map_codegen_31.cpp @@ -45,11 +45,11 @@ // CK31A-USE: [[MTYPE00:@.+]] = private {{.*}}constant [7 x i64] [i64 [[#0x1020]], // CK31A-NOUSE: [[MTYPE00:@.+]] = private {{.*}}constant [7 x i64] [i64 [[#0x1000]], // -// MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003 // MEMBER_OF_1=0x1000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x1000000001003 +// MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003 // PRESENT=0x1000 | TARGET_PARAM=0x20 | FROM=0x2 | TO=0x1 = 0x1023 -// CK31A-USE-SAME: {{^}} i64 [[#0x1000000000003]], i64 [[#0x1000000001003]], i64 [[#0x1023]], -// CK31A-NOUSE-SAME: {{^}} i64 [[#0x1000000000003]], i64 [[#0x1000000001003]], i64 [[#0x1003]], +// CK31A-USE-SAME: {{^}} i64 [[#0x1000000001003]], i64 [[#0x1000000000003]], i64 [[#0x1023]], +// CK31A-NOUSE-SAME: {{^}} i64 [[#0x1000000001003]], i64 [[#0x1000000000003]], i64 [[#0x1003]], // // PRESENT=0x1000 | TARGET_PARAM=0x20 = 0x1020 // MEMBER_OF_5=0x5000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x5000000001003 @@ -77,91 +77,94 @@ struct ST st1; struct ST st2; - // Make sure the struct picks up present even if another element of the struct - // doesn't have present. - // Region 00 - // CK31A: [[ST1_I:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST1]], i{{.+}} 0, i{{.+}} 0 - // CK31A: [[ST1_J:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST1]], i{{.+}} 0, i{{.+}} 1 - // CK31A: [[ST2_I:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST2]], i{{.+}} 0, i{{.+}} 0 - // CK31A: [[ST2_J:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST2]], i{{.+}} 0, i{{.+}} 1 - // CK31A-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{[^,]+}}, i8* {{[^,]+}}, i32 7, i8** [[GEPBP:%[0-9]+]], i8** [[GEPP:%[0-9]+]], i64* [[GEPS:%.+]], i64* getelementptr {{.+}}[7 x i{{.+}}]* [[MTYPE00]]{{.+}}) - // CK31A-DAG: [[GEPS]] = getelementptr inbounds [7 x i64], [7 x i64]* [[S:%.+]], i{{.+}} 0, i{{.+}} 0 - // CK31A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] - // CK31A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] - - // st1 - // CK31A-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 - // CK31A-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 - // CK31A-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 - // CK31A-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]** - // CK31A-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** - // CK31A-DAG: store [[ST]]* [[ST1]], [[ST]]** [[CBP0]] - // CK31A-DAG: store i32* [[ST1_I]], i32** [[CP0]] - // CK31A-DAG: store i64 %{{.+}}, i64* [[S0]] - - // st1.i - // CK31A-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 - // CK31A-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 - // CK31A-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 - // CK31A-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]** - // CK31A-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32** - // CK31A-DAG: store [[ST]]* [[ST1]], [[ST]]** [[CBP1]] - // CK31A-DAG: store i32* [[ST1_I]], i32** [[CP1]] - // CK31A-DAG: store i64 4, i64* [[S1]] - - // st1.j - // CK31A-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 - // CK31A-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 - // CK31A-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2 - // CK31A-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to [[ST]]** - // CK31A-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to i32** - // CK31A-DAG: store [[ST]]* [[ST1]], [[ST]]** [[CBP2]] - // CK31A-DAG: store i32* [[ST1_J]], i32** [[CP2]] - // CK31A-DAG: store i64 4, i64* [[S2]] - - // a - // CK31A-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3 - // CK31A-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3 - // CK31A-DAG: [[S3:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 3 - // CK31A-DAG: [[CBP3:%.+]] = bitcast i8** [[BP3]] to i32** - // CK31A-DAG: [[CP3:%.+]] = bitcast i8** [[P3]] to i32** - // CK31A-DAG: store i32* [[A]], i32** [[CBP3]] - // CK31A-DAG: store i32* [[A]], i32** [[CP3]] - // CK31A-DAG: store i64 4, i64* [[S3]] - - // st2 - // CK31A-DAG: [[BP4:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 4 - // CK31A-DAG: [[P4:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 4 - // CK31A-DAG: [[S4:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 4 - // CK31A-DAG: [[CBP4:%.+]] = bitcast i8** [[BP4]] to [[ST]]** - // CK31A-DAG: [[CP4:%.+]] = bitcast i8** [[P4]] to i32** - // CK31A-DAG: store [[ST]]* [[ST2]], [[ST]]** [[CBP4]] - // CK31A-DAG: store i32* [[ST2_I]], i32** [[CP4]] - // CK31A-DAG: store i64 %{{.+}}, i64* [[S4]] - - // st2.i - // CK31A-DAG: [[BP5:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 5 - // CK31A-DAG: [[P5:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 5 - // CK31A-DAG: [[S5:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 5 - // CK31A-DAG: [[CBP5:%.+]] = bitcast i8** [[BP5]] to [[ST]]** - // CK31A-DAG: [[CP5:%.+]] = bitcast i8** [[P5]] to i32** - // CK31A-DAG: store [[ST]]* [[ST2]], [[ST]]** [[CBP5]] - // CK31A-DAG: store i32* [[ST2_I]], i32** [[CP5]] - // CK31A-DAG: store i64 4, i64* [[S5]] - - // st2.j - // CK31A-DAG: [[BP6:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 6 - // CK31A-DAG: [[P6:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 6 - // CK31A-DAG: [[S6:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 6 - // CK31A-DAG: [[CBP6:%.+]] = bitcast i8** [[BP6]] to [[ST]]** - // CK31A-DAG: [[CP6:%.+]] = bitcast i8** [[P6]] to i32** - // CK31A-DAG: store [[ST]]* [[ST2]], [[ST]]** [[CBP6]] - // CK31A-DAG: store i32* [[ST2_J]], i32** [[CP6]] - // CK31A-DAG: store i64 4, i64* [[S6]] - - // CK31A-USE: call void [[CALL00:@.+]]([[ST]]* [[ST1]], i32* [[A]], [[ST]]* [[ST2]]) - // CK31A-NOUSE: call void [[CALL00:@.+]]() - #pragma omp target map(tofrom: st1.i) map(present, tofrom: a, st1.j, st2.i) map(tofrom: st2.j) +// Make sure the struct picks up present even if another element of the struct +// doesn't have present. +// Region 00 +// CK31A: [[ST1_J:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST1]], i{{.+}} 0, i{{.+}} 1 +// CK31A: [[ST1_I:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST1]], i{{.+}} 0, i{{.+}} 0 +// CK31A: [[ST2_I:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST2]], i{{.+}} 0, i{{.+}} 0 +// CK31A: [[ST2_J:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST2]], i{{.+}} 0, i{{.+}} 1 +// CK31A-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{[^,]+}}, i8* {{[^,]+}}, i32 7, i8** [[GEPBP:%[0-9]+]], i8** [[GEPP:%[0-9]+]], i64* [[GEPS:%.+]], i64* getelementptr {{.+}}[7 x i{{.+}}]* [[MTYPE00]]{{.+}}) +// CK31A-DAG: [[GEPS]] = getelementptr inbounds [7 x i64], [7 x i64]* [[S:%.+]], i{{.+}} 0, i{{.+}} 0 +// CK31A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] +// CK31A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + +// st1 +// CK31A-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 +// CK31A-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 +// CK31A-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 +// CK31A-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]** +// CK31A-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** +// CK31A-DAG: store [[ST]]* [[ST1]], [[ST]]** [[CBP0]] +// CK31A-DAG: store i32* [[ST1_I]], i32** [[CP0]] +// CK31A-DAG: store i64 %{{.+}}, i64* [[S0]] + +// st1.j +// CK31A-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 +// CK31A-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 +// CK31A-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 +// CK31A-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]** +// CK31A-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32** +// CK31A-DAG: store [[ST]]* [[ST1]], [[ST]]** [[CBP1]] +// CK31A-DAG: store i32* [[ST1_J]], i32** [[CP1]] +// CK31A-DAG: store i64 4, i64* [[S1]] + +// st1.i +// CK31A-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 +// CK31A-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 +// CK31A-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2 +// CK31A-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to [[ST]]** +// CK31A-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to i32** +// CK31A-DAG: store [[ST]]* [[ST1]], [[ST]]** [[CBP2]] +// CK31A-DAG: store i32* [[ST1_I]], i32** [[CP2]] +// CK31A-DAG: store i64 4, i64* [[S2]] + +// a +// CK31A-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3 +// CK31A-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3 +// CK31A-DAG: [[S3:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 3 +// CK31A-DAG: [[CBP3:%.+]] = bitcast i8** [[BP3]] to i32** +// CK31A-DAG: [[CP3:%.+]] = bitcast i8** [[P3]] to i32** +// CK31A-DAG: store i32* [[A]], i32** [[CBP3]] +// CK31A-DAG: store i32* [[A]], i32** [[CP3]] +// CK31A-DAG: store i64 4, i64* [[S3]] + +// st2 +// CK31A-DAG: [[BP4:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 4 +// CK31A-DAG: [[P4:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 4 +// CK31A-DAG: [[S4:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 4 +// CK31A-DAG: [[CBP4:%.+]] = bitcast i8** [[BP4]] to [[ST]]** +// CK31A-DAG: [[CP4:%.+]] = bitcast i8** [[P4]] to i32** +// CK31A-DAG: store [[ST]]* [[ST2]], [[ST]]** [[CBP4]] +// CK31A-DAG: store i32* [[ST2_I]], i32** [[CP4]] +// CK31A-DAG: store i64 %{{.+}}, i64* [[S4]] + +// st2.i +// CK31A-DAG: [[BP5:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 5 +// CK31A-DAG: [[P5:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 5 +// CK31A-DAG: [[S5:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 5 +// CK31A-DAG: [[CBP5:%.+]] = bitcast i8** [[BP5]] to [[ST]]** +// CK31A-DAG: [[CP5:%.+]] = bitcast i8** [[P5]] to i32** +// CK31A-DAG: store [[ST]]* [[ST2]], [[ST]]** [[CBP5]] +// CK31A-DAG: store i32* [[ST2_I]], i32** [[CP5]] +// CK31A-DAG: store i64 4, i64* [[S5]] + +// st2.j +// CK31A-DAG: [[BP6:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 6 +// CK31A-DAG: [[P6:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 6 +// CK31A-DAG: [[S6:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 6 +// CK31A-DAG: [[CBP6:%.+]] = bitcast i8** [[BP6]] to [[ST]]** +// CK31A-DAG: [[CP6:%.+]] = bitcast i8** [[P6]] to i32** +// CK31A-DAG: store [[ST]]* [[ST2]], [[ST]]** [[CBP6]] +// CK31A-DAG: store i32* [[ST2_J]], i32** [[CP6]] +// CK31A-DAG: store i64 4, i64* [[S6]] + +// CK31A-USE: call void [[CALL00:@.+]]([[ST]]* [[ST1]], i32* [[A]], [[ST]]* [[ST2]]) +// CK31A-NOUSE: call void [[CALL00:@.+]]() +#pragma omp target map(tofrom \ + : st1.i) map(present, tofrom \ + : a, st1.j, st2.i) map(tofrom \ + : st2.j) { #ifdef USE st1.i++; diff --git a/clang/test/OpenMP/target_map_codegen_32.cpp b/clang/test/OpenMP/target_map_codegen_32.cpp --- a/clang/test/OpenMP/target_map_codegen_32.cpp +++ b/clang/test/OpenMP/target_map_codegen_32.cpp @@ -36,63 +36,65 @@ // CK31B: [[ST:%.+]] = type { i32, i32 } // PRESENT=0x1000 | TARGET_PARAM=0x20 = 0x1020 -// MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003 // MEMBER_OF_1=0x1000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x1000000001003 +// MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003 // CK31B-LABEL: @.__omp_offloading_{{.*}}test_present_members{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 // CK31B-USE: [[MTYPE00:@.+]] = private {{.*}}constant [3 x i64] [i64 [[#0x1020]], // CK31B-NOUSE: [[MTYPE00:@.+]] = private {{.*}}constant [3 x i64] [i64 [[#0x1000]], -// CK31B-USE-SAME: {{^}} i64 [[#0x1000000000003]], i64 [[#0x1000000001003]]] -// CK31B-NOUSE-SAME: {{^}} i64 [[#0x1000000000003]], i64 [[#0x1000000001003]]] +// CK31B-USE-SAME: {{^}} i64 [[#0x1000000001003]], i64 [[#0x1000000000003]]] +// CK31B-NOUSE-SAME: {{^}} i64 [[#0x1000000001003]], i64 [[#0x1000000000003]]] struct ST { int i; int j; // CK31B-LABEL: define {{.*}}test_present_members{{.*}}( void test_present_members() { - // Make sure the struct picks up present even if another element of the - // struct doesn't have present. - // Region 00 - // CK31B: [[I:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[THIS:%.+]], i{{.+}} 0, i{{.+}} 0 - // CK31B: [[J:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[THIS]], i{{.+}} 0, i{{.+}} 1 - // CK31B-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{[^,]+}}, i8* {{[^,]+}}, i32 3, i8** [[GEPBP:%[0-9]+]], i8** [[GEPP:%[0-9]+]], i64* [[GEPS:%.+]], i64* getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE00]]{{.+}}) - // CK31B-DAG: [[GEPS]] = getelementptr inbounds [3 x i64], [3 x i64]* [[S:%.+]], i{{.+}} 0, i{{.+}} 0 - // CK31B-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] - // CK31B-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] +// Make sure the struct picks up present even if another element of the +// struct doesn't have present. +// Region 00 +// CK31B: [[J:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[THIS:%.+]], i{{.+}} 0, i{{.+}} 1 +// CK31B: [[I:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[THIS]], i{{.+}} 0, i{{.+}} 0 +// CK31B-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{[^,]+}}, i8* {{[^,]+}}, i32 3, i8** [[GEPBP:%[0-9]+]], i8** [[GEPP:%[0-9]+]], i64* [[GEPS:%.+]], i64* getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE00]]{{.+}}) +// CK31B-DAG: [[GEPS]] = getelementptr inbounds [3 x i64], [3 x i64]* [[S:%.+]], i{{.+}} 0, i{{.+}} 0 +// CK31B-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] +// CK31B-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] - // this - // CK31B-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 - // CK31B-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 - // CK31B-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 - // CK31B-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]** - // CK31B-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** - // CK31B-DAG: store [[ST]]* [[THIS]], [[ST]]** [[CBP0]] - // CK31B-DAG: store i32* [[I]], i32** [[CP0]] - // CK31B-DAG: store i64 %{{.+}}, i64* [[S0]] +// this +// CK31B-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 +// CK31B-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 +// CK31B-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 +// CK31B-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]** +// CK31B-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** +// CK31B-DAG: store [[ST]]* [[THIS]], [[ST]]** [[CBP0]] +// CK31B-DAG: store i32* [[I]], i32** [[CP0]] +// CK31B-DAG: store i64 %{{.+}}, i64* [[S0]] - // i - // CK31B-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 - // CK31B-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 - // CK31B-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 - // CK31B-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]** - // CK31B-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32** - // CK31B-DAG: store [[ST]]* [[THIS]], [[ST]]** [[CBP1]] - // CK31B-DAG: store i32* [[I]], i32** [[CP1]] - // CK31B-DAG: store i64 4, i64* [[S1]] +// j +// CK31B-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 +// CK31B-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 +// CK31B-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 +// CK31B-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]** +// CK31B-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32** +// CK31B-DAG: store [[ST]]* [[THIS]], [[ST]]** [[CBP1]] +// CK31B-DAG: store i32* [[J]], i32** [[CP1]] +// CK31B-DAG: store i64 4, i64* [[S1]] - // j - // CK31B-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 - // CK31B-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 - // CK31B-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2 - // CK31B-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to [[ST]]** - // CK31B-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to i32** - // CK31B-DAG: store [[ST]]* [[THIS]], [[ST]]** [[CBP2]] - // CK31B-DAG: store i32* [[J]], i32** [[CP2]] - // CK31B-DAG: store i64 4, i64* [[S2]] +// i +// CK31B-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 +// CK31B-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 +// CK31B-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2 +// CK31B-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to [[ST]]** +// CK31B-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to i32** +// CK31B-DAG: store [[ST]]* [[THIS]], [[ST]]** [[CBP2]] +// CK31B-DAG: store i32* [[I]], i32** [[CP2]] +// CK31B-DAG: store i64 4, i64* [[S2]] - // CK31B-USE: call void [[CALL00:@.+]]([[ST]]* [[THIS]]) - // CK31B-NOUSE: call void [[CALL00:@.+]]() - #pragma omp target map(tofrom: i) map(present, tofrom: j) +// CK31B-USE: call void [[CALL00:@.+]]([[ST]]* [[THIS]]) +// CK31B-NOUSE: call void [[CALL00:@.+]]() +#pragma omp target map(tofrom \ + : i) map(present, tofrom \ + : j) { #ifdef USE i++; diff --git a/clang/test/OpenMP/target_map_messages.cpp b/clang/test/OpenMP/target_map_messages.cpp --- a/clang/test/OpenMP/target_map_messages.cpp +++ b/clang/test/OpenMP/target_map_messages.cpp @@ -79,7 +79,7 @@ #pragma omp target map(to:b,e) {} - #pragma omp target map(to:b,e) map(to:b) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} + #pragma omp target map(to:b,e) map(to:b) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}} {} #pragma omp target map(to:b[:2],e) {} @@ -338,19 +338,19 @@ {} #pragma omp target map(r.ArrS[0].Error) // expected-error {{no member named 'Error' in 'SB'}} {} -#pragma omp target map(r.ArrS[0].A, r.ArrS[1].A) // expected-error {{multiple array elements associated with the same variable are not allowed in map clauses of the same construct}} expected-note {{used here}} +#pragma omp target map(r.ArrS[0].A, r.ArrS[1].A) // lt50-error {{multiple array elements associated with the same variable are not allowed in map clauses of the same construct}} lt50-note {{used here}} {} #pragma omp target map(r.ArrS[0].A, t.ArrS[1].A) {} -#pragma omp target map(r.PtrS[0], r.PtrS->B) // expected-error {{same pointer dereferenced in multiple different ways in map clause expressions}} expected-note {{used here}} +#pragma omp target map(r.PtrS[0], r.PtrS->B) // lt50-error {{same pointer dereferenced in multiple different ways in map clause expressions}} lt50-note {{used here}} {} -#pragma omp target map(r.PtrS, r.PtrS->B) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}} +#pragma omp target map(r.PtrS, r.PtrS->B) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}} {} #pragma omp target map(r.PtrS->A, r.PtrS->B) {} -#pragma omp target map(r.RPtrS[0], r.RPtrS->B) // expected-error {{same pointer dereferenced in multiple different ways in map clause expressions}} expected-note {{used here}} +#pragma omp target map(r.RPtrS[0], r.RPtrS->B) // lt50-error {{same pointer dereferenced in multiple different ways in map clause expressions}} lt50-note {{used here}} {} -#pragma omp target map(r.RPtrS, r.RPtrS->B) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}} +#pragma omp target map(r.RPtrS, r.RPtrS->B) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}} {} #pragma omp target map(r.RPtrS->A, r.RPtrS->B) {} @@ -360,13 +360,13 @@ {} #pragma omp target map(r.C, r.D) {} -#pragma omp target map(r.C, r.C) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} +#pragma omp target map(r.C, r.C) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}} {} -#pragma omp target map(r.C) map(r.C) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} +#pragma omp target map(r.C) map(r.C) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}} {} #pragma omp target map(r.C, r.S) // this would be an error only caught at runtime - Sema would have to make sure there is not way for the missing data between fields to be mapped somewhere else. {} -#pragma omp target map(r, r.S) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} +#pragma omp target map(r, r.S) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}} {} #pragma omp target map(r.C, t.C) {} @@ -401,16 +401,16 @@ } #pragma omp target data map(to \ - : r.C) //expected-note {{used here}} + : r.C) // lt50-note {{used here}} { -#pragma omp target map(r.D) // expected-error {{original storage of expression in data environment is shared but data environment do not fully contain mapped expression storage}} +#pragma omp target map(r.D) // lt50-error {{original storage of expression in data environment is shared but data environment do not fully contain mapped expression storage}} {} } #pragma omp target data map(to \ - : t.Ptr) //expected-note {{used here}} + : t.Ptr) // lt50-note {{used here}} { -#pragma omp target map(t.Ptr[:23]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target map(t.Ptr[:23]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} {} } @@ -581,19 +581,19 @@ #pragma omp target data map(S2::S2sc) #pragma omp target data map(e, g) // warn-warning 2 {{Type 'S4' is not trivially copyable and not guaranteed to be mapped correctly}} warn-warning 2 {{Type 'S5' is not trivially copyable and not guaranteed to be mapped correctly}} #pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}} -#pragma omp target data map(k) map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}} -#pragma omp target map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}} +#pragma omp target data map(k) map(k) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}} +#pragma omp target map(k), map(k[:5]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} lt50-note 2 {{used here}} foo(); #pragma omp target data map(da) #pragma omp target map(da[:4]) foo(); -#pragma omp target data map(k, j, l) // expected-note 2 {{used here}} -#pragma omp target data map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target data map(k, j, l) // lt50-note 2 {{used here}} +#pragma omp target data map(k[:4]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} #pragma omp target data map(j) -#pragma omp target map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}} +#pragma omp target map(l) map(l[:5]) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}} foo(); -#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}} -#pragma omp target data map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target data map(k[:4], j, l[:5]) // lt50-note 2 {{used here}} +#pragma omp target data map(k) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} #pragma omp target data map(j) #pragma omp target map(l) foo(); @@ -716,19 +716,19 @@ #pragma omp target data map(S2::S2sc) #pragma omp target data map(e, g) // warn-warning {{Type 'S4' is not trivially copyable and not guaranteed to be mapped correctly}} warn-warning {{Type 'S5' is not trivially copyable and not guaranteed to be mapped correctly}} #pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}} -#pragma omp target data map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} -#pragma omp target map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}} +#pragma omp target data map(k), map(k) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}} +#pragma omp target map(k), map(k[:5]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}} foo(); #pragma omp target data map(da) #pragma omp target map(da[:4]) foo(); -#pragma omp target data map(k, j, l) // expected-note {{used here}} -#pragma omp target data map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target data map(k, j, l) // lt50-note {{used here}} +#pragma omp target data map(k[:4]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} #pragma omp target data map(j) -#pragma omp target map(l) map(l[:5]) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} +#pragma omp target map(l) map(l[:5]) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}} foo(); -#pragma omp target data map(k[:4], j, l[:5]) // expected-note {{used here}} -#pragma omp target data map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target data map(k[:4], j, l[:5]) // lt50-note {{used here}} +#pragma omp target data map(k) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} #pragma omp target data map(j) #pragma omp target map(l) foo(); @@ -766,47 +766,49 @@ {} #pragma omp target map(m) // warn-warning {{Type 'S6' is not trivially copyable and not guaranteed to be mapped correctly}} {} -// expected-note@+1 {{used here}} +#pragma omp target + { s.a++; } +// lt50-note@+1 {{used here}} #pragma omp target map(s.s.s) -// expected-error@+1 {{variable already marked as mapped in current construct}} +// lt50-error@+1 {{variable already marked as mapped in current construct}} { s.a++; } -// expected-note@+1 {{used here}} +// lt50-note@+1 {{used here}} #pragma omp target map(s.s.s.a) -// expected-error@+1 {{variable already marked as mapped in current construct}} +// lt50-error@+1 {{variable already marked as mapped in current construct}} { s.a++; } -// expected-note@+1 {{used here}} +// lt50-note@+1 {{used here}} #pragma omp target map(s.b[:5]) -// expected-error@+1 {{variable already marked as mapped in current construct}} +// lt50-error@+1 {{variable already marked as mapped in current construct}} { s.a++; } #pragma omp target map(s.p[:5]) { s.a++; } -// expected-note@+1 {{used here}} +// lt50-note@+1 {{used here}} #pragma omp target map(s.s.sa[3].a) -// expected-error@+1 {{variable already marked as mapped in current construct}} +// lt50-error@+1 {{variable already marked as mapped in current construct}} { s.a++; } -// expected-note@+1 {{used here}} +// lt50-note@+1 {{used here}} #pragma omp target map(s.s.sp[3]->a) -// expected-error@+1 {{variable already marked as mapped in current construct}} +// lt50-error@+1 {{variable already marked as mapped in current construct}} { s.a++; } -// expected-note@+1 {{used here}} +// lt50-note@+1 {{used here}} #pragma omp target map(s.p->a) -// expected-error@+1 {{variable already marked as mapped in current construct}} +// lt50-error@+1 {{variable already marked as mapped in current construct}} { s.a++; } -// expected-note@+1 {{used here}} +// lt50-note@+1 {{used here}} #pragma omp target map(s.s.p->a) -// expected-error@+1 {{variable already marked as mapped in current construct}} +// lt50-error@+1 {{variable already marked as mapped in current construct}} { s.a++; } -// expected-note@+1 {{used here}} +// lt50-note@+1 {{used here}} #pragma omp target map(s.s.s.b[:2]) -// expected-error@+1 {{variable already marked as mapped in current construct}} +// lt50-error@+1 {{variable already marked as mapped in current construct}} { s.a++; } -// expected-note@+1 {{used here}} +// lt50-note@+1 {{used here}} #pragma omp target map(s.s.p->b[:2]) -// expected-error@+1 {{variable already marked as mapped in current construct}} +// lt50-error@+1 {{variable already marked as mapped in current construct}} { s.a++; } -// expected-note@+1 {{used here}} +// lt50-note@+1 {{used here}} #pragma omp target map(s.p->p->p->a) -// expected-error@+1 {{variable already marked as mapped in current construct}} +// lt50-error@+1 {{variable already marked as mapped in current construct}} { s.a++; } #pragma omp target map(s.s.s.b[:2]) { s.s.s.b[0]++; } diff --git a/clang/test/OpenMP/target_parallel_for_map_messages.cpp b/clang/test/OpenMP/target_parallel_for_map_messages.cpp --- a/clang/test/OpenMP/target_parallel_for_map_messages.cpp +++ b/clang/test/OpenMP/target_parallel_for_map_messages.cpp @@ -152,25 +152,26 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}} for (i = 0; i < argc; ++i) foo(); -#pragma omp target parallel for map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}} +#pragma omp target parallel for map(k), map(k) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}} for (i = 0; i < argc; ++i) foo(); -#pragma omp target parallel for map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}} +#pragma omp target parallel for map(k), map(k[:5]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} lt50-note 2 {{used here}} for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for map(da) for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for map(da[:4]) for (i = 0; i < argc; ++i) foo(); -#pragma omp target data map(k, j, l) // expected-note 2 {{used here}} -#pragma omp target parallel for map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target data map(k, j, l) // lt50-note 2 {{used here}} +#pragma omp target parallel for map(k[:4]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for map(j) for (i = 0; i < argc; ++i) foo(); -#pragma omp target parallel for map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}} - for (i = 0; i < argc; ++i) foo(); -#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}} -{ -#pragma omp target parallel for map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target parallel for map(l) map(l[:5]) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}} for (i = 0; i < argc; ++i) foo(); +#pragma omp target data map(k[:4], j, l[:5]) // lt50-note 2 {{used here}} + { +#pragma omp target parallel for map(k) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} + for (i = 0; i < argc; ++i) + foo(); #pragma omp target parallel for map(j) for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for map(l) @@ -269,25 +270,26 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}} for (i = 0; i < argc; ++i) foo(); -#pragma omp target parallel for map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} +#pragma omp target parallel for map(k), map(k) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}} for (i = 0; i < argc; ++i) foo(); -#pragma omp target parallel for map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}} +#pragma omp target parallel for map(k), map(k[:5]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}} for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for map(da) for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for map(da[:4]) for (i = 0; i < argc; ++i) foo(); -#pragma omp target data map(k, j, l) // expected-note {{used here}} -#pragma omp target parallel for map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target data map(k, j, l) // lt50-note {{used here}} +#pragma omp target parallel for map(k[:4]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for map(j) for (i = 0; i < argc; ++i) foo(); -#pragma omp target parallel for map(l) map(l[:5]) // expected-error 1 {{variable already marked as mapped in current construct}} expected-note 1 {{used here}} - for (i = 0; i < argc; ++i) foo(); -#pragma omp target data map(k[:4], j, l[:5]) // expected-note {{used here}} -{ -#pragma omp target parallel for map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target parallel for map(l) map(l[:5]) // lt50-error 1 {{variable already marked as mapped in current construct}} lt50-note 1 {{used here}} for (i = 0; i < argc; ++i) foo(); +#pragma omp target data map(k[:4], j, l[:5]) // lt50-note {{used here}} + { +#pragma omp target parallel for map(k) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} + for (i = 0; i < argc; ++i) + foo(); #pragma omp target parallel for map(j) for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for map(l) diff --git a/clang/test/OpenMP/target_parallel_for_simd_map_messages.cpp b/clang/test/OpenMP/target_parallel_for_simd_map_messages.cpp --- a/clang/test/OpenMP/target_parallel_for_simd_map_messages.cpp +++ b/clang/test/OpenMP/target_parallel_for_simd_map_messages.cpp @@ -152,25 +152,26 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for simd map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}} for (i = 0; i < argc; ++i) foo(); -#pragma omp target parallel for simd map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}} +#pragma omp target parallel for simd map(k), map(k) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}} for (i = 0; i < argc; ++i) foo(); -#pragma omp target parallel for simd map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}} +#pragma omp target parallel for simd map(k), map(k[:5]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} lt50-note 2 {{used here}} for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for simd map(da) for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for simd map(da[:4]) for (i = 0; i < argc; ++i) foo(); -#pragma omp target data map(k, j, l) // expected-note 2 {{used here}} -#pragma omp target parallel for simd map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target data map(k, j, l) // lt50-note 2 {{used here}} +#pragma omp target parallel for simd map(k[:4]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for simd map(j) for (i = 0; i < argc; ++i) foo(); -#pragma omp target parallel for simd map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}} - for (i = 0; i < argc; ++i) foo(); -#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}} -{ -#pragma omp target parallel for simd map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target parallel for simd map(l) map(l[:5]) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}} for (i = 0; i < argc; ++i) foo(); +#pragma omp target data map(k[:4], j, l[:5]) // lt50-note 2 {{used here}} + { +#pragma omp target parallel for simd map(k) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} + for (i = 0; i < argc; ++i) + foo(); #pragma omp target parallel for simd map(j) for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for simd map(l) @@ -269,25 +270,26 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for simd map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}} for (i = 0; i < argc; ++i) foo(); -#pragma omp target parallel for simd map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} +#pragma omp target parallel for simd map(k), map(k) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}} for (i = 0; i < argc; ++i) foo(); -#pragma omp target parallel for simd map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}} +#pragma omp target parallel for simd map(k), map(k[:5]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}} for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for simd map(da) for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for simd map(da[:4]) for (i = 0; i < argc; ++i) foo(); -#pragma omp target data map(k, j, l) // expected-note {{used here}} -#pragma omp target parallel for simd map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target data map(k, j, l) // lt50-note {{used here}} +#pragma omp target parallel for simd map(k[:4]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for simd map(j) for (i = 0; i < argc; ++i) foo(); -#pragma omp target parallel for simd map(l) map(l[:5]) // expected-error 1 {{variable already marked as mapped in current construct}} expected-note 1 {{used here}} - for (i = 0; i < argc; ++i) foo(); -#pragma omp target data map(k[:4], j, l[:5]) // expected-note {{used here}} -{ -#pragma omp target parallel for simd map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target parallel for simd map(l) map(l[:5]) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}} for (i = 0; i < argc; ++i) foo(); +#pragma omp target data map(k[:4], j, l[:5]) // lt50-note {{used here}} + { +#pragma omp target parallel for simd map(k) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} + for (i = 0; i < argc; ++i) + foo(); #pragma omp target parallel for simd map(j) for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for simd map(l) diff --git a/clang/test/OpenMP/target_parallel_map_messages.cpp b/clang/test/OpenMP/target_parallel_map_messages.cpp --- a/clang/test/OpenMP/target_parallel_map_messages.cpp +++ b/clang/test/OpenMP/target_parallel_map_messages.cpp @@ -152,25 +152,25 @@ foo(); #pragma omp target parallel map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}} foo(); -#pragma omp target parallel map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}} +#pragma omp target parallel map(k), map(k) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}} foo(); -#pragma omp target parallel map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}} +#pragma omp target parallel map(k), map(k[:5]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} lt50-note 2 {{used here}} foo(); #pragma omp target parallel map(da) foo(); #pragma omp target parallel map(da[:4]) foo(); -#pragma omp target data map(k, j, l) // expected-note 2 {{used here}} -#pragma omp target parallel map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target data map(k, j, l) // lt50-note 2 {{used here}} +#pragma omp target parallel map(k[:4]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} foo(); #pragma omp target parallel map(j) foo(); -#pragma omp target parallel map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}} - foo(); -#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}} -{ -#pragma omp target parallel map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target parallel map(l) map(l[:5]) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}} foo(); +#pragma omp target data map(k[:4], j, l[:5]) // lt50-note 2 {{used here}} + { +#pragma omp target parallel map(k) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} + foo(); #pragma omp target parallel map(j) foo(); #pragma omp target parallel map(l) @@ -267,25 +267,25 @@ foo(); #pragma omp target parallel map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}} foo(); -#pragma omp target parallel map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} +#pragma omp target parallel map(k), map(k) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}} foo(); -#pragma omp target parallel map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}} +#pragma omp target parallel map(k), map(k[:5]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}} foo(); #pragma omp target parallel map(da) foo(); #pragma omp target parallel map(da[:4]) foo(); -#pragma omp target data map(k, j, l) // expected-note {{used here}} -#pragma omp target parallel map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target data map(k, j, l) // lt50-note {{used here}} +#pragma omp target parallel map(k[:4]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} foo(); #pragma omp target parallel map(j) foo(); -#pragma omp target parallel map(l) map(l[:5]) // expected-error 1 {{variable already marked as mapped in current construct}} expected-note 1 {{used here}} - foo(); -#pragma omp target data map(k[:4], j, l[:5]) // expected-note 1 {{used here}} -{ -#pragma omp target parallel map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target parallel map(l) map(l[:5]) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}} foo(); +#pragma omp target data map(k[:4], j, l[:5]) // lt50-note {{used here}} + { +#pragma omp target parallel map(k) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} + foo(); #pragma omp target parallel map(j) foo(); #pragma omp target parallel map(l) diff --git a/clang/test/OpenMP/target_simd_map_messages.cpp b/clang/test/OpenMP/target_simd_map_messages.cpp --- a/clang/test/OpenMP/target_simd_map_messages.cpp +++ b/clang/test/OpenMP/target_simd_map_messages.cpp @@ -146,25 +146,26 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target simd map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}} for (i = 0; i < argc; ++i) foo(); -#pragma omp target simd map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}} +#pragma omp target simd map(k), map(k) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}} for (i = 0; i < argc; ++i) foo(); -#pragma omp target simd map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}} +#pragma omp target simd map(k), map(k[:5]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} lt50-note 2 {{used here}} for (i = 0; i < argc; ++i) foo(); #pragma omp target simd map(da) for (i = 0; i < argc; ++i) foo(); #pragma omp target simd map(da[:4]) for (i = 0; i < argc; ++i) foo(); -#pragma omp target data map(k, j, l) // expected-note 2 {{used here}} -#pragma omp target simd map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target data map(k, j, l) // lt50-note 2 {{used here}} +#pragma omp target simd map(k[:4]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} for (i = 0; i < argc; ++i) foo(); #pragma omp target simd map(j) for (i = 0; i < argc; ++i) foo(); -#pragma omp target simd map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}} - for (i = 0; i < argc; ++i) foo(); -#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}} -{ -#pragma omp target simd map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target simd map(l) map(l[:5]) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}} for (i = 0; i < argc; ++i) foo(); +#pragma omp target data map(k[:4], j, l[:5]) // lt50-note 2 {{used here}} + { +#pragma omp target simd map(k) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} + for (i = 0; i < argc; ++i) + foo(); #pragma omp target simd map(j) for (i = 0; i < argc; ++i) foo(); #pragma omp target simd map(l) // OK @@ -256,31 +257,37 @@ #pragma omp target simd map(e, g) for (i = 0; i < argc; ++i) foo(); #pragma omp target simd map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}} - for (i = 0; i < argc; ++i) foo(); -#pragma omp target simd map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} - for (i = 0; i < argc; ++i) foo(); -#pragma omp target simd map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}} + for (i = 0; i < argc; ++i) + foo(); +#pragma omp target simd map(k), map(k) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}} + for (i = 0; i < argc; ++i) + foo(); +#pragma omp target simd map(k), map(k[:5]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}} for (i = 0; i < argc; ++i) foo(); #pragma omp target simd map(da) for (i = 0; i < argc; ++i) foo(); #pragma omp target simd map(da[:4]) - for (i = 0; i < argc; ++i) foo(); -#pragma omp target data map(k, j, l) // expected-note {{used here}} -#pragma omp target simd map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} + for (i = 0; i < argc; ++i) + foo(); +#pragma omp target data map(k, j, l) // lt50-note {{used here}} +#pragma omp target simd map(k[:4]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} for (i = 0; i < argc; ++i) foo(); #pragma omp target simd map(j) - for (i = 0; i < argc; ++i) foo(); -#pragma omp target simd map(l) map(l[:5]) // expected-error 1 {{variable already marked as mapped in current construct}} expected-note 1 {{used here}} - for (i = 0; i < argc; ++i) foo(); -#pragma omp target data map(k[:4], j, l[:5]) // expected-note 1 {{used here}} -{ -#pragma omp target simd map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}} - for (i = 0; i < argc; ++i) foo(); + for (i = 0; i < argc; ++i) + foo(); +#pragma omp target simd map(l) map(l[:5]) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}} + for (i = 0; i < argc; ++i) + foo(); +#pragma omp target data map(k[:4], j, l[:5]) // lt50-note {{used here}} + { +#pragma omp target simd map(k) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} + for (i = 0; i < argc; ++i) + foo(); #pragma omp target simd map(j) for (i = 0; i < argc; ++i) foo(); #pragma omp target simd map(l) // for (i = 0; i < argc; ++i) foo(); -} + } #pragma omp target simd map(always, tofrom: x) for (i = 0; i < argc; ++i) foo(); diff --git a/clang/test/OpenMP/target_teams_distribute_map_messages.cpp b/clang/test/OpenMP/target_teams_distribute_map_messages.cpp --- a/clang/test/OpenMP/target_teams_distribute_map_messages.cpp +++ b/clang/test/OpenMP/target_teams_distribute_map_messages.cpp @@ -152,25 +152,26 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}} for (i = 0; i < argc; ++i) foo(); -#pragma omp target teams distribute map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}} +#pragma omp target teams distribute map(k), map(k) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}} for (i = 0; i < argc; ++i) foo(); -#pragma omp target teams distribute map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}} +#pragma omp target teams distribute map(k), map(k[:5]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} lt50-note 2 {{used here}} for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute map(da) for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute map(da[:4]) for (i = 0; i < argc; ++i) foo(); -#pragma omp target data map(k, j, l) // expected-note 2 {{used here}} -#pragma omp target teams distribute map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target data map(k, j, l) // lt50-note 2 {{used here}} +#pragma omp target teams distribute map(k[:4]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute map(j) for (i = 0; i < argc; ++i) foo(); -#pragma omp target teams distribute map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}} - for (i = 0; i < argc; ++i) foo(); -#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}} -{ -#pragma omp target teams distribute map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target teams distribute map(l) map(l[:5]) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}} for (i = 0; i < argc; ++i) foo(); +#pragma omp target data map(k[:4], j, l[:5]) // lt50-note 2 {{used here}} + { +#pragma omp target teams distribute map(k) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} + for (i = 0; i < argc; ++i) + foo(); #pragma omp target teams distribute map(j) for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute map(l) @@ -269,25 +270,26 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}} for (i = 0; i < argc; ++i) foo(); -#pragma omp target teams distribute map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} +#pragma omp target teams distribute map(k), map(k) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}} for (i = 0; i < argc; ++i) foo(); -#pragma omp target teams distribute map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}} +#pragma omp target teams distribute map(k), map(k[:5]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}} for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute map(da) for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute map(da[:4]) for (i = 0; i < argc; ++i) foo(); -#pragma omp target data map(k, j, l) // expected-note {{used here}} -#pragma omp target teams distribute map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target data map(k, j, l) // lt50-note {{used here}} +#pragma omp target teams distribute map(k[:4]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute map(j) for (i = 0; i < argc; ++i) foo(); -#pragma omp target teams distribute map(l) map(l[:5]) // expected-error 1 {{variable already marked as mapped in current construct}} expected-note 1 {{used here}} - for (i = 0; i < argc; ++i) foo(); -#pragma omp target data map(k[:4], j, l[:5]) // expected-note {{used here}} -{ -#pragma omp target teams distribute map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target teams distribute map(l) map(l[:5]) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}} for (i = 0; i < argc; ++i) foo(); +#pragma omp target data map(k[:4], j, l[:5]) // lt50-note {{used here}} + { +#pragma omp target teams distribute map(k) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} + for (i = 0; i < argc; ++i) + foo(); #pragma omp target teams distribute map(j) for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute map(l) diff --git a/clang/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp b/clang/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp --- a/clang/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp +++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp @@ -150,25 +150,26 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}} for (i = 0; i < argc; ++i) foo(); -#pragma omp target teams distribute parallel for map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}} +#pragma omp target teams distribute parallel for map(k), map(k) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}} for (i = 0; i < argc; ++i) foo(); -#pragma omp target teams distribute parallel for map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}} +#pragma omp target teams distribute parallel for map(k), map(k[:5]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} lt50-note 2 {{used here}} for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for map(da) for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for map(da[:4]) for (i = 0; i < argc; ++i) foo(); -#pragma omp target data map(k, j, l) // expected-note 2 {{used here}} -#pragma omp target teams distribute parallel for map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target data map(k, j, l) // lt50-note 2 {{used here}} +#pragma omp target teams distribute parallel for map(k[:4]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for map(j) for (i = 0; i < argc; ++i) foo(); -#pragma omp target teams distribute parallel for map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}} - for (i = 0; i < argc; ++i) foo(); -#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}} -{ -#pragma omp target teams distribute parallel for map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target teams distribute parallel for map(l) map(l[:5]) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}} for (i = 0; i < argc; ++i) foo(); +#pragma omp target data map(k[:4], j, l[:5]) // lt50-note 2 {{used here}} + { +#pragma omp target teams distribute parallel for map(k) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} + for (i = 0; i < argc; ++i) + foo(); #pragma omp target teams distribute parallel for map(j) for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for map(l) @@ -265,31 +266,37 @@ #pragma omp target teams distribute parallel for map(e, g) for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}} - for (i = 0; i < argc; ++i) foo(); -#pragma omp target teams distribute parallel for map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} - for (i = 0; i < argc; ++i) foo(); -#pragma omp target teams distribute parallel for map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}} + for (i = 0; i < argc; ++i) + foo(); +#pragma omp target teams distribute parallel for map(k), map(k) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}} + for (i = 0; i < argc; ++i) + foo(); +#pragma omp target teams distribute parallel for map(k), map(k[:5]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}} for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for map(da) for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for map(da[:4]) - for (i = 0; i < argc; ++i) foo(); -#pragma omp target data map(k, j, l) // expected-note {{used here}} -#pragma omp target teams distribute parallel for map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} + for (i = 0; i < argc; ++i) + foo(); +#pragma omp target data map(k, j, l) // lt50-note {{used here}} +#pragma omp target teams distribute parallel for map(k[:4]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for map(j) - for (i = 0; i < argc; ++i) foo(); -#pragma omp target teams distribute parallel for map(l) map(l[:5]) // expected-error 1 {{variable already marked as mapped in current construct}} expected-note 1 {{used here}} - for (i = 0; i < argc; ++i) foo(); -#pragma omp target data map(k[:4], j, l[:5]) // expected-note {{used here}} -{ -#pragma omp target teams distribute parallel for map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}} - for (i = 0; i < argc; ++i) foo(); + for (i = 0; i < argc; ++i) + foo(); +#pragma omp target teams distribute parallel for map(l) map(l[:5]) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}} + for (i = 0; i < argc; ++i) + foo(); +#pragma omp target data map(k[:4], j, l[:5]) // lt50-note {{used here}} + { +#pragma omp target teams distribute parallel for map(k) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} + for (i = 0; i < argc; ++i) + foo(); #pragma omp target teams distribute parallel for map(j) for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for map(l) for (i = 0; i < argc; ++i) foo(); -} + } #pragma omp target teams distribute parallel for map(always, tofrom: x) for (i = 0; i < argc; ++i) foo(); diff --git a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp --- a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp +++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp @@ -152,25 +152,26 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for simd map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}} for (i = 0; i < argc; ++i) foo(); -#pragma omp target teams distribute parallel for simd map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}} +#pragma omp target teams distribute parallel for simd map(k), map(k) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}} for (i = 0; i < argc; ++i) foo(); -#pragma omp target teams distribute parallel for simd map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}} +#pragma omp target teams distribute parallel for simd map(k), map(k[:5]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} lt50-note 2 {{used here}} for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for simd map(da) for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for simd map(da[:4]) for (i = 0; i < argc; ++i) foo(); -#pragma omp target data map(k, j, l) // expected-note 2 {{used here}} -#pragma omp target teams distribute parallel for simd map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target data map(k, j, l) // lt50-note 2 {{used here}} +#pragma omp target teams distribute parallel for simd map(k[:4]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for simd map(j) for (i = 0; i < argc; ++i) foo(); -#pragma omp target teams distribute parallel for simd map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}} - for (i = 0; i < argc; ++i) foo(); -#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}} -{ -#pragma omp target teams distribute parallel for simd map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target teams distribute parallel for simd map(l) map(l[:5]) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}} for (i = 0; i < argc; ++i) foo(); +#pragma omp target data map(k[:4], j, l[:5]) // lt50-note 2 {{used here}} + { +#pragma omp target teams distribute parallel for simd map(k) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} + for (i = 0; i < argc; ++i) + foo(); #pragma omp target teams distribute parallel for simd map(j) for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for simd map(l) @@ -268,25 +269,26 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for simd map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}} for (i = 0; i < argc; ++i) foo(); -#pragma omp target teams distribute parallel for simd map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} +#pragma omp target teams distribute parallel for simd map(k), map(k) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}} for (i = 0; i < argc; ++i) foo(); -#pragma omp target teams distribute parallel for simd map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}} +#pragma omp target teams distribute parallel for simd map(k), map(k[:5]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}} for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for simd map(da) for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for simd map(da[:4]) for (i = 0; i < argc; ++i) foo(); -#pragma omp target data map(k, j, l) // expected-note {{used here}} -#pragma omp target teams distribute parallel for simd map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target data map(k, j, l) // lt50-note {{used here}} +#pragma omp target teams distribute parallel for simd map(k[:4]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for simd map(j) for (i = 0; i < argc; ++i) foo(); -#pragma omp target teams distribute parallel for simd map(l) map(l[:5]) // expected-error 1 {{variable already marked as mapped in current construct}} expected-note 1 {{used here}} - for (i = 0; i < argc; ++i) foo(); -#pragma omp target data map(k[:4], j, l[:5]) // expected-note {{used here}} -{ -#pragma omp target teams distribute parallel for simd map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target teams distribute parallel for simd map(l) map(l[:5]) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}} for (i = 0; i < argc; ++i) foo(); +#pragma omp target data map(k[:4], j, l[:5]) // lt50-note {{used here}} + { +#pragma omp target teams distribute parallel for simd map(k) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} + for (i = 0; i < argc; ++i) + foo(); #pragma omp target teams distribute parallel for simd map(j) for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute parallel for simd map(l) diff --git a/clang/test/OpenMP/target_teams_distribute_simd_map_messages.cpp b/clang/test/OpenMP/target_teams_distribute_simd_map_messages.cpp --- a/clang/test/OpenMP/target_teams_distribute_simd_map_messages.cpp +++ b/clang/test/OpenMP/target_teams_distribute_simd_map_messages.cpp @@ -152,25 +152,26 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute simd map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}} for (i = 0; i < argc; ++i) foo(); -#pragma omp target teams distribute simd map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}} +#pragma omp target teams distribute simd map(k), map(k) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}} for (i = 0; i < argc; ++i) foo(); -#pragma omp target teams distribute simd map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}} +#pragma omp target teams distribute simd map(k), map(k[:5]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} lt50-note 2 {{used here}} for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute simd map(da) for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute simd map(da[:4]) for (i = 0; i < argc; ++i) foo(); -#pragma omp target data map(k, j, l) // expected-note 2 {{used here}} -#pragma omp target teams distribute simd map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target data map(k, j, l) // lt50-note 2 {{used here}} +#pragma omp target teams distribute simd map(k[:4]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute simd map(j) for (i = 0; i < argc; ++i) foo(); -#pragma omp target teams distribute simd map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}} - for (i = 0; i < argc; ++i) foo(); -#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}} -{ -#pragma omp target teams distribute simd map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target teams distribute simd map(l) map(l[:5]) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}} for (i = 0; i < argc; ++i) foo(); +#pragma omp target data map(k[:4], j, l[:5]) // lt50-note 2 {{used here}} + { +#pragma omp target teams distribute simd map(k) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} + for (i = 0; i < argc; ++i) + foo(); #pragma omp target teams distribute simd map(j) for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute simd map(l) @@ -269,25 +270,26 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute simd map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}} for (i = 0; i < argc; ++i) foo(); -#pragma omp target teams distribute simd map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} +#pragma omp target teams distribute simd map(k), map(k) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}} for (i = 0; i < argc; ++i) foo(); -#pragma omp target teams distribute simd map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}} +#pragma omp target teams distribute simd map(k), map(k[:5]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}} for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute simd map(da) for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute simd map(da[:4]) for (i = 0; i < argc; ++i) foo(); -#pragma omp target data map(k, j, l) // expected-note {{used here}} -#pragma omp target teams distribute simd map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target data map(k, j, l) // lt50-note {{used here}} +#pragma omp target teams distribute simd map(k[:4]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute simd map(j) for (i = 0; i < argc; ++i) foo(); -#pragma omp target teams distribute simd map(l) map(l[:5]) // expected-error 1 {{variable already marked as mapped in current construct}} expected-note 1 {{used here}} - for (i = 0; i < argc; ++i) foo(); -#pragma omp target data map(k[:4], j, l[:5]) // expected-note {{used here}} -{ -#pragma omp target teams distribute simd map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target teams distribute simd map(l) map(l[:5]) // lt50-error 1 {{variable already marked as mapped in current construct}} lt50-note 1 {{used here}} for (i = 0; i < argc; ++i) foo(); +#pragma omp target data map(k[:4], j, l[:5]) // lt50-note {{used here}} + { +#pragma omp target teams distribute simd map(k) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} + for (i = 0; i < argc; ++i) + foo(); #pragma omp target teams distribute simd map(j) for (i = 0; i < argc; ++i) foo(); #pragma omp target teams distribute simd map(l) diff --git a/clang/test/OpenMP/target_teams_map_messages.cpp b/clang/test/OpenMP/target_teams_map_messages.cpp --- a/clang/test/OpenMP/target_teams_map_messages.cpp +++ b/clang/test/OpenMP/target_teams_map_messages.cpp @@ -62,7 +62,7 @@ #pragma omp target teams map(to:b,e) {} - #pragma omp target teams map(to:b,e) map(to:b) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} + #pragma omp target teams map(to:b,e) map(to:b) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}} {} #pragma omp target teams map(to:b[:2],e) {} @@ -238,19 +238,19 @@ {} #pragma omp target teams map(r.ArrS[0].Error) // expected-error {{no member named 'Error' in 'SB'}} {} - #pragma omp target teams map(r.ArrS[0].A, r.ArrS[1].A) // expected-error {{multiple array elements associated with the same variable are not allowed in map clauses of the same construct}} expected-note {{used here}} + #pragma omp target teams map(r.ArrS[0].A, r.ArrS[1].A) // lt50-error {{multiple array elements associated with the same variable are not allowed in map clauses of the same construct}} lt50-note {{used here}} {} #pragma omp target teams map(r.ArrS[0].A, t.ArrS[1].A) {} - #pragma omp target teams map(r.PtrS[0], r.PtrS->B) // expected-error {{same pointer dereferenced in multiple different ways in map clause expressions}} expected-note {{used here}} + #pragma omp target teams map(r.PtrS[0], r.PtrS->B) // lt50-error {{same pointer dereferenced in multiple different ways in map clause expressions}} lt50-note {{used here}} {} - #pragma omp target teams map(r.PtrS, r.PtrS->B) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}} + #pragma omp target teams map(r.PtrS, r.PtrS->B) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}} {} #pragma omp target teams map(r.PtrS->A, r.PtrS->B) {} - #pragma omp target teams map(r.RPtrS[0], r.RPtrS->B) // expected-error {{same pointer dereferenced in multiple different ways in map clause expressions}} expected-note {{used here}} + #pragma omp target teams map(r.RPtrS[0], r.RPtrS->B) // lt50-error {{same pointer dereferenced in multiple different ways in map clause expressions}} lt50-note {{used here}} {} - #pragma omp target teams map(r.RPtrS, r.RPtrS->B) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}} + #pragma omp target teams map(r.RPtrS, r.RPtrS->B) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}} {} #pragma omp target teams map(r.RPtrS->A, r.RPtrS->B) {} @@ -262,13 +262,13 @@ {} #pragma omp target teams map(r.C, r.D) {} - #pragma omp target teams map(r.C, r.C) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} + #pragma omp target teams map(r.C, r.C) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}} {} - #pragma omp target teams map(r.C) map(r.C) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} + #pragma omp target teams map(r.C) map(r.C) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}} {} #pragma omp target teams map(r.C, r.S) // this would be an error only caught at runtime - Sema would have to make sure there is not way for the missing data between fields to be mapped somewhere else. {} - #pragma omp target teams map(r, r.S) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} + #pragma omp target teams map(r, r.S) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}} {} #pragma omp target teams map(r.C, t.C) {} @@ -298,15 +298,15 @@ #pragma omp target teams map(u.B) // expected-error {{mapping of union members is not allowed}} {} - #pragma omp target data map(to: r.C) //expected-note {{used here}} + #pragma omp target data map(to: r.C) // lt50-note {{used here}} { - #pragma omp target teams map(r.D) // expected-error {{original storage of expression in data environment is shared but data environment do not fully contain mapped expression storage}} + #pragma omp target teams map(r.D) // lt50-error {{original storage of expression in data environment is shared but data environment do not fully contain mapped expression storage}} {} } - #pragma omp target data map(to: t.Ptr) //expected-note {{used here}} + #pragma omp target data map(to: t.Ptr) // lt50-note {{used here}} { - #pragma omp target teams map(t.Ptr[:23]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} + #pragma omp target teams map(t.Ptr[:23]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} {} } @@ -457,22 +457,22 @@ #pragma omp target data map(S2::S2sc) #pragma omp target data map(e, g) #pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}} -#pragma omp target data map(k) map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}} -#pragma omp target teams map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}} +#pragma omp target data map(k) map(k) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}} +#pragma omp target teams map(k), map(k[:5]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} lt50-note 2 {{used here}} foo(); #pragma omp target data map(da) #pragma omp target teams map(da[:4]) foo(); -#pragma omp target data map(k, j, l) // expected-note 2 {{used here}} -#pragma omp target data map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target data map(k, j, l) // lt50-note 2 {{used here}} +#pragma omp target data map(k[:4]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} #pragma omp target data map(j) -#pragma omp target teams map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}} +#pragma omp target teams map(l) map(l[:5]) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}} foo(); -#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}} -#pragma omp target data map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target data map(k[:4], j, l[:5]) // lt50-note 2 {{used here}} +#pragma omp target data map(k) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} #pragma omp target data map(j) #pragma omp target teams map(l) foo(); @@ -540,22 +540,22 @@ #pragma omp target data map(S2::S2sc) #pragma omp target data map(e, g) #pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}} -#pragma omp target data map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} -#pragma omp target teams map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}} +#pragma omp target data map(k), map(k) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}} +#pragma omp target teams map(k), map(k[:5]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}} foo(); #pragma omp target data map(da) #pragma omp target teams map(da[:4]) foo(); -#pragma omp target data map(k, j, l) // expected-note {{used here}} -#pragma omp target data map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target data map(k, j, l) // lt50-note {{used here}} +#pragma omp target data map(k[:4]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} #pragma omp target data map(j) -#pragma omp target teams map(l) map(l[:5]) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} +#pragma omp target teams map(l) map(l[:5]) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}} foo(); -#pragma omp target data map(k[:4], j, l[:5]) // expected-note {{used here}} -#pragma omp target data map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}} +#pragma omp target data map(k[:4], j, l[:5]) // lt50-note {{used here}} +#pragma omp target data map(k) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} #pragma omp target data map(j) #pragma omp target teams map(l) foo(); diff --git a/clang/test/OpenMP/target_update_codegen.cpp b/clang/test/OpenMP/target_update_codegen.cpp --- a/clang/test/OpenMP/target_update_codegen.cpp +++ b/clang/test/OpenMP/target_update_codegen.cpp @@ -1113,7 +1113,7 @@ // SIMD-ONLY0-NOT: {{__kmpc|__tgt}} #ifdef CK19 -// PRESENT=0x1000 | TO=0x1 = 0x1021 +// PRESENT=0x1000 | TO=0x1 = 0x1001 // CK19: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1001]]] // PRESENT=0x1000 | FROM=0x2 = 0x1002 diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp --- a/openmp/libomptarget/src/omptarget.cpp +++ b/openmp/libomptarget/src/omptarget.cpp @@ -258,7 +258,7 @@ for (unsigned I = 0, E = MapperComponents.Components.size(); I < E; ++I) { auto &C = MapperComponents - .Components[target_data_function == targetDataEnd ? I : E - I - 1]; + .Components[target_data_function == targetDataEnd ? E - I - 1 : I]; MapperArgsBase[I] = C.Base; MapperArgs[I] = C.Begin; MapperArgSizes[I] = C.Size; diff --git a/openmp/libomptarget/test/mapping/declare_mapper_nested_mappers.cpp b/openmp/libomptarget/test/mapping/declare_mapper_nested_mappers.cpp new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/mapping/declare_mapper_nested_mappers.cpp @@ -0,0 +1,66 @@ +// RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu +// RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu +// RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu +// RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu +// RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda + +#include +#include + +typedef struct { + int a; + double *b; +} C; +#pragma omp declare mapper(id1 : C s) map(to : s.a) map(from : s.b [0:2]) + +typedef struct { + int e; + C f; + int h; + short *g; +} D; +#pragma omp declare mapper(default \ + : D r) map(from \ + : r.e) map(mapper(id1), tofrom \ + : r.f) map(tofrom \ + : r.g [0:r.h]) + +int main() { + constexpr int N = 10; + D s; + s.e = 111; + s.f.a = 222; + double x[2]; + x[1] = 20; + short y[N]; + y[1] = 30; + s.f.b = &x[0]; + s.g = &y[0]; + s.h = N; + + D *sp = &s; + D **spp = &sp; + + printf("%d %d %4.5f %d %d %d\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.b[1], + spp[0][0].f.b == &x[0] ? 1 : 0, spp[0][0].g[1], + spp[0][0].g == &y[0] ? 1 : 0); + // CHECK: 111 222 20.00000 1 30 1 + + __intptr_t p = reinterpret_cast<__intptr_t>(&x[0]), + p1 = reinterpret_cast<__intptr_t>(&y[0]); +#pragma omp target map(tofrom : spp[0][0]) firstprivate(p, p1) + { + printf("%d %d %d %d\n", spp[0][0].f.a, + spp[0][0].f.b == reinterpret_cast(p) ? 1 : 0, spp[0][0].g[1], + spp[0][0].g == reinterpret_cast(p1) ? 1 : 0); + // CHECK: 222 0 30 0 + spp[0][0].e = 333; + spp[0][0].f.a = 444; + spp[0][0].f.b[1] = 40; + spp[0][0].g[1] = 50; + } + printf("%d %d %4.5f %d %d %d\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.b[1], + spp[0][0].f.b == &x[0] ? 1 : 0, spp[0][0].g[1], + spp[0][0].g == &y[0] ? 1 : 0); + // CHECK: 333 222 40.00000 1 50 1 +}