[clang] 60d71a2 - [OPENMP50]Allow overlapping mapping in target constructs.
Alexey Bataev via cfe-commits
cfe-commits at lists.llvm.org
Tue Feb 16 14:47:45 PST 2021
Author: Alexey Bataev
Date: 2021-02-16T14:42:08-08:00
New Revision: 60d71a286b5a03653fc99cd09423d603feb897de
URL: https://github.com/llvm/llvm-project/commit/60d71a286b5a03653fc99cd09423d603feb897de
DIFF: https://github.com/llvm/llvm-project/commit/60d71a286b5a03653fc99cd09423d603feb897de.diff
LOG: [OPENMP50]Allow overlapping mapping in target constructs.
OpenMP 5.0 removed a lot of restriction for overlapped mapped items
comparing to OpenMP 4.5. Patch restricts the checks for overlapped data
mappings only for OpenMP 4.5 and less and reorders mapping of the
arguments so, that present and alloc mappings are processed first and
then all others.
Differential Revision: https://reviews.llvm.org/D86119
Added:
openmp/libomptarget/test/mapping/declare_mapper_nested_mappers.cpp
Modified:
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/lib/Sema/SemaOpenMP.cpp
clang/test/OpenMP/declare_mapper_codegen.cpp
clang/test/OpenMP/target_data_codegen.cpp
clang/test/OpenMP/target_enter_data_codegen.cpp
clang/test/OpenMP/target_is_device_ptr_messages.cpp
clang/test/OpenMP/target_map_codegen_29.cpp
clang/test/OpenMP/target_map_codegen_31.cpp
clang/test/OpenMP/target_map_codegen_32.cpp
clang/test/OpenMP/target_map_messages.cpp
clang/test/OpenMP/target_parallel_for_map_messages.cpp
clang/test/OpenMP/target_parallel_for_simd_map_messages.cpp
clang/test/OpenMP/target_parallel_map_messages.cpp
clang/test/OpenMP/target_simd_map_messages.cpp
clang/test/OpenMP/target_teams_distribute_map_messages.cpp
clang/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp
clang/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp
clang/test/OpenMP/target_teams_distribute_simd_map_messages.cpp
clang/test/OpenMP/target_teams_map_messages.cpp
clang/test/OpenMP/target_update_codegen.cpp
openmp/libomptarget/src/omptarget.cpp
Removed:
################################################################################
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 3a1949d15b87..a56dad819ac0 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7157,6 +7157,7 @@ class MappableExprsHandler {
0, Address::invalid()};
Address Base = Address::invalid();
bool IsArraySection = false;
+ bool HasCompleteRecord = false;
};
private:
@@ -7339,11 +7340,10 @@ class MappableExprsHandler {
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 @@ class MappableExprsHandler {
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 @@ class MappableExprsHandler {
}
}
-public:
- MappableExprsHandler(const OMPExecutableDirective &Dir, CodeGenFunction &CGF)
- : CurDir(&Dir), CGF(CGF) {
- // Extract firstprivate clause information.
- for (const auto *C : Dir.getClausesOfKind<OMPFirstprivateClause>())
- for (const auto *D : C->varlists())
- FirstPrivateDecls.try_emplace(
- cast<VarDecl>(cast<DeclRefExpr>(D)->getDecl()), C->isImplicit());
- // Extract implicit firstprivates from uses_allocators clauses.
- for (const auto *C : Dir.getClausesOfKind<OMPUsesAllocatorsClause>()) {
- for (unsigned I = 0, E = C->getNumberOfAllocators(); I < E; ++I) {
- OMPUsesAllocatorsClause::Data D = C->getAllocatorData(I);
- if (const auto *DRE = dyn_cast_or_null<DeclRefExpr>(D.AllocatorTraits))
- FirstPrivateDecls.try_emplace(cast<VarDecl>(DRE->getDecl()),
- /*Implicit=*/true);
- else if (const auto *VD = dyn_cast<VarDecl>(
- cast<DeclRefExpr>(D.Allocator->IgnoreParenImpCasts())
- ->getDecl()))
- FirstPrivateDecls.try_emplace(VD, /*Implicit=*/true);
- }
- }
- // Extract device pointer clause information.
- for (const auto *C : Dir.getClausesOfKind<OMPIsDevicePtrClause>())
- 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<const OMPClause *> Clauses, MapCombinedInfoTy &CombinedInfo,
const llvm::DenseSet<CanonicalDeclPtr<const Decl>> &SkipVarSet =
llvm::DenseSet<CanonicalDeclPtr<const Decl>>()) 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<const ValueDecl *, SmallVector<MapInfo, 8>> Info;
+ enum MapKind { Present, Allocs, Other, Total };
+ llvm::MapVector<CanonicalDeclPtr<const Decl>,
+ SmallVector<SmallVector<MapInfo, 8>, 4>>
+ Info;
// Helper function to fill the information map for the
diff erent supported
// clauses.
auto &&InfoGen =
[&Info, &SkipVarSet](
- const ValueDecl *D,
+ const ValueDecl *D, MapKind Kind,
OMPClauseMappableExprCommon::MappableExprComponentListRef L,
OpenMPMapClauseKind MapType,
ArrayRef<OpenMPMapModifierKind> MapModifiers,
ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
bool ReturnDevicePointer, bool IsImplicit, const ValueDecl *Mapper,
const Expr *VarRef = nullptr, bool ForDeviceAddr = false) {
- const ValueDecl *VD =
- D ? cast<ValueDecl>(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<SmallVector<MapInfo, 8>, 4>(Total)))
+ .first;
+ It->second[Kind].emplace_back(
+ L, MapType, MapModifiers, MotionModifiers, ReturnDevicePointer,
+ IsImplicit, Mapper, VarRef, ForDeviceAddr);
};
- assert(CurDir.is<const OMPExecutableDirective *>() &&
- "Expect a executable directive");
- const auto *CurExecDir = CurDir.get<const OMPExecutableDirective *>();
- for (const auto *C : CurExecDir->getClausesOfKind<OMPMapClause>()) {
+ for (const auto *Cl : Clauses) {
+ const auto *C = dyn_cast<OMPMapClause>(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<OMPToClause>()) {
+ for (const auto *Cl : Clauses) {
+ const auto *C = dyn_cast<OMPToClause>(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<OMPFromClause>()) {
+ for (const auto *Cl : Clauses) {
+ const auto *C = dyn_cast<OMPFromClause>(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 @@ class MappableExprsHandler {
// 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<const ValueDecl *, SmallVector<DeferredDevicePtrEntryTy, 4>>
+ llvm::MapVector<CanonicalDeclPtr<const Decl>,
+ SmallVector<DeferredDevicePtrEntryTy, 4>>
DeferredInfo;
MapCombinedInfoTy UseDevicePtrCombinedInfo;
- for (const auto *C :
- CurExecDir->getClausesOfKind<OMPUseDevicePtrClause>()) {
+ for (const auto *Cl : Clauses) {
+ const auto *C = dyn_cast<OMPUseDevicePtrClause>(Cl);
+ if (!C)
+ continue;
for (const auto L : C->component_lists()) {
OMPClauseMappableExprCommon::MappableExprComponentListRef Components =
std::get<1>(L);
@@ -8384,28 +8340,34 @@ class MappableExprsHandler {
// 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<VarDecl>(VD);
- if (CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory() ||
- isa<MemberExpr>(IE) ||
- !VD->getType().getNonReferenceType()->isPointerType() ||
- PrevCI == CI->Components.rend() ||
- isa<MemberExpr>(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<VarDecl>(VD);
+ if (CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory() ||
+ isa<MemberExpr>(IE) ||
+ !VD->getType().getNonReferenceType()->isPointerType() ||
+ PrevCI == CI->Components.rend() ||
+ isa<MemberExpr>(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 @@ class MappableExprsHandler {
// 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 @@ class MappableExprsHandler {
// 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<CanonicalDeclPtr<const Decl>, 4> Processed;
- for (const auto *C :
- CurExecDir->getClausesOfKind<OMPUseDeviceAddrClause>()) {
+ for (const auto *Cl : Clauses) {
+ const auto *C = dyn_cast<OMPUseDeviceAddrClause>(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 @@ class MappableExprsHandler {
// 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 @@ class MappableExprsHandler {
// 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 @@ class MappableExprsHandler {
}
}
- 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<ValueDecl>(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 @@ class MappableExprsHandler {
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 @@ class MappableExprsHandler {
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<const OMPDeclareMapperDecl *>() &&
- "Expect a declare mapper directive");
- const auto *CurMapperDir = CurDir.get<const OMPDeclareMapperDecl *>();
- // 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<const ValueDecl *, SmallVector<MapInfo, 8>> Info;
-
- // Fill the information map for map clauses.
- for (const auto *C : CurMapperDir->clauselists()) {
- const auto *MC = cast<OMPMapClause>(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<ValueDecl>(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<OMPFirstprivateClause>())
+ for (const auto *D : C->varlists())
+ FirstPrivateDecls.try_emplace(
+ cast<VarDecl>(cast<DeclRefExpr>(D)->getDecl()), C->isImplicit());
+ // Extract implicit firstprivates from uses_allocators clauses.
+ for (const auto *C : Dir.getClausesOfKind<OMPUsesAllocatorsClause>()) {
+ for (unsigned I = 0, E = C->getNumberOfAllocators(); I < E; ++I) {
+ OMPUsesAllocatorsClause::Data D = C->getAllocatorData(I);
+ if (const auto *DRE = dyn_cast_or_null<DeclRefExpr>(D.AllocatorTraits))
+ FirstPrivateDecls.try_emplace(cast<VarDecl>(DRE->getDecl()),
+ /*Implicit=*/true);
+ else if (const auto *VD = dyn_cast<VarDecl>(
+ cast<DeclRefExpr>(D.Allocator->IgnoreParenImpCasts())
+ ->getDecl()))
+ FirstPrivateDecls.try_emplace(VD, /*Implicit=*/true);
}
}
+ // Extract device pointer clause information.
+ for (const auto *C : Dir.getClausesOfKind<OMPIsDevicePtrClause>())
+ 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<CanonicalDeclPtr<const Decl>> &SkipVarSet =
+ llvm::DenseSet<CanonicalDeclPtr<const Decl>>()) const {
+ assert(CurDir.is<const OMPExecutableDirective *>() &&
+ "Expect a executable directive");
+ const auto *CurExecDir = CurDir.get<const OMPExecutableDirective *>();
+ 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<const OMPDeclareMapperDecl *>() &&
+ "Expect a declare mapper directive");
+ const auto *CurMapperDir = CurDir.get<const OMPDeclareMapperDecl *>();
+ generateAllInfoForClauses(CurMapperDir->clauses(), CombinedInfo);
}
/// Emit capture info for lambdas for variables captured by reference.
@@ -8806,6 +8827,25 @@ class MappableExprsHandler {
++EI;
}
}
+ llvm::stable_sort(DeclComponentLists, [](const MapData &LHS,
+ const MapData &RHS) {
+ ArrayRef<OpenMPMapModifierKind> 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 @@ class MappableExprsHandler {
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<MemberExpr>(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 @@ class MappableExprsHandler {
}
}
for (auto &Pair : OverlappedData) {
- llvm::sort(
+ llvm::stable_sort(
Pair.getSecond(),
[&Layout](
OMPClauseMappableExprCommon::MappableExprComponentListRef First,
@@ -8908,6 +8960,7 @@ class MappableExprsHandler {
// 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 @@ class MappableExprsHandler {
L;
ArrayRef<OMPClauseMappableExprCommon::MappableExprComponentListRef>
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 @@ getNestedDistributeDirective(ASTContext &Ctx, const OMPExecutableDirective &D) {
/// 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 @@ getNestedDistributeDirective(ASTContext &Ctx, const OMPExecutableDirective &D) {
/// }
/// }
/// // 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 @@ void CGOpenMPRuntime::emitUserDefinedMapper(const OMPDeclareMapperDecl *D,
// 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 @@ void CGOpenMPRuntime::emitUserDefinedMapper(const OMPDeclareMapperDecl *D,
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 @@ void CGOpenMPRuntime::emitUserDefinedMapper(const OMPDeclareMapperDecl *D,
.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<PointerType>())
- .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 @@ void CGOpenMPRuntime::emitUserDefinedMapper(const OMPDeclareMapperDecl *D,
: 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 @@ void CGOpenMPRuntime::emitUDMapperArrayInitOrDel(
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 @@ void CGOpenMPRuntime::emitUDMapperArrayInitOrDel(
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 @@ void CGOpenMPRuntime::emitTargetCall(
llvm::DenseSet<CanonicalDeclPtr<const Decl>> 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
index 3791665197d9..bc7e8b55c757 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -3564,9 +3564,11 @@ class DSAAttrChecker final : public StmtVisitor<DSAAttrChecker, void> {
!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 @@ static bool checkMapConflicts(
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 @@ static void checkMappableExpressionList(
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
index 4bb9ca975dc2..8d820e8e3355 100644
--- a/clang/test/OpenMP/declare_mapper_codegen.cpp
+++ b/clang/test/OpenMP/declare_mapper_codegen.cpp
@@ -97,17 +97,21 @@ class C {
// 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 @@ class C {
// 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 @@ class C {
// 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 @@ class C {
// 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 @@ class C {
// 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 @@ class C {
// 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 @@ class C {
// 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 @@ class C {
// 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 @@ class C {
// 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 @@ class C {
// 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 @@ class C {
// 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 @@ class C {
// 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 @@ class C {
// 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 @@ class C {
// 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 @@ class C {
// 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 @@ class C {
// 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
index 9728d3ed3518..2666a736dc58 100644
--- a/clang/test/OpenMP/target_data_codegen.cpp
+++ b/clang/test/OpenMP/target_data_codegen.cpp
@@ -561,8 +561,9 @@ struct S2 {
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 @@ void test_close_modifier(int arg) {
// 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 @@ void test_present_modifier(int arg) {
// 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 @@ void test_present_modifier(int arg) {
// 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
index c87fb0aa795b..8d9eaa352c5c 100644
--- a/clang/test/OpenMP/target_enter_data_codegen.cpp
+++ b/clang/test/OpenMP/target_enter_data_codegen.cpp
@@ -267,7 +267,7 @@ double gc[100];
// 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
index 92297438d525..7137f5deff19 100644
--- 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
index 59b0476a76c5..360a44812197 100644
--- a/clang/test/OpenMP/target_map_codegen_29.cpp
+++ b/clang/test/OpenMP/target_map_codegen_29.cpp
@@ -73,9 +73,8 @@ typedef struct StructWithPtrTag : public Base {
// 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 @@ typedef struct StructWithPtrTag : public Base {
// 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
index d1874a8eae18..6e4e9dd0a6a9 100644
--- a/clang/test/OpenMP/target_map_codegen_31.cpp
+++ b/clang/test/OpenMP/target_map_codegen_31.cpp
@@ -45,11 +45,11 @@ struct ST {
// 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 @@ void explicit_maps_single (int ii){
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
index 4c6c19c061b1..8f1b6c5bbada 100644
--- 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
index 7311f6acd823..7f39b3cbd118 100644
--- a/clang/test/OpenMP/target_map_messages.cpp
+++ b/clang/test/OpenMP/target_map_messages.cpp
@@ -79,7 +79,7 @@ struct SA {
#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 @@ void SAclient(int arg) {
{}
#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
diff erent 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
diff erent 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
diff erent 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
diff erent 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 @@ void SAclient(int arg) {
{}
#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 @@ void SAclient(int arg) {
}
#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 @@ T tmain(T argc) {
#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 @@ int main(int argc, char **argv) {
#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 @@ int main(int argc, char **argv) {
{}
#pragma omp target map(m) // warn-warning {{Type 'S6<int>' is not trivially copyable and not guaranteed to be mapped correctly}}
{}
-// expected-note at +1 {{used here}}
+#pragma omp target
+ { s.a++; }
+// lt50-note at +1 {{used here}}
#pragma omp target map(s.s.s)
-// expected-error at +1 {{variable already marked as mapped in current construct}}
+// lt50-error at +1 {{variable already marked as mapped in current construct}}
{ s.a++; }
-// expected-note at +1 {{used here}}
+// lt50-note at +1 {{used here}}
#pragma omp target map(s.s.s.a)
-// expected-error at +1 {{variable already marked as mapped in current construct}}
+// lt50-error at +1 {{variable already marked as mapped in current construct}}
{ s.a++; }
-// expected-note at +1 {{used here}}
+// lt50-note at +1 {{used here}}
#pragma omp target map(s.b[:5])
-// expected-error at +1 {{variable already marked as mapped in current construct}}
+// lt50-error at +1 {{variable already marked as mapped in current construct}}
{ s.a++; }
#pragma omp target map(s.p[:5])
{ s.a++; }
-// expected-note at +1 {{used here}}
+// lt50-note at +1 {{used here}}
#pragma omp target map(s.s.sa[3].a)
-// expected-error at +1 {{variable already marked as mapped in current construct}}
+// lt50-error at +1 {{variable already marked as mapped in current construct}}
{ s.a++; }
-// expected-note at +1 {{used here}}
+// lt50-note at +1 {{used here}}
#pragma omp target map(s.s.sp[3]->a)
-// expected-error at +1 {{variable already marked as mapped in current construct}}
+// lt50-error at +1 {{variable already marked as mapped in current construct}}
{ s.a++; }
-// expected-note at +1 {{used here}}
+// lt50-note at +1 {{used here}}
#pragma omp target map(s.p->a)
-// expected-error at +1 {{variable already marked as mapped in current construct}}
+// lt50-error at +1 {{variable already marked as mapped in current construct}}
{ s.a++; }
-// expected-note at +1 {{used here}}
+// lt50-note at +1 {{used here}}
#pragma omp target map(s.s.p->a)
-// expected-error at +1 {{variable already marked as mapped in current construct}}
+// lt50-error at +1 {{variable already marked as mapped in current construct}}
{ s.a++; }
-// expected-note at +1 {{used here}}
+// lt50-note at +1 {{used here}}
#pragma omp target map(s.s.s.b[:2])
-// expected-error at +1 {{variable already marked as mapped in current construct}}
+// lt50-error at +1 {{variable already marked as mapped in current construct}}
{ s.a++; }
-// expected-note at +1 {{used here}}
+// lt50-note at +1 {{used here}}
#pragma omp target map(s.s.p->b[:2])
-// expected-error at +1 {{variable already marked as mapped in current construct}}
+// lt50-error at +1 {{variable already marked as mapped in current construct}}
{ s.a++; }
-// expected-note at +1 {{used here}}
+// lt50-note at +1 {{used here}}
#pragma omp target map(s.p->p->p->a)
-// expected-error at +1 {{variable already marked as mapped in current construct}}
+// lt50-error at +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
index 665ecfa926aa..a1ebb2d0adfd 100644
--- a/clang/test/OpenMP/target_parallel_for_map_messages.cpp
+++ b/clang/test/OpenMP/target_parallel_for_map_messages.cpp
@@ -152,25 +152,26 @@ T tmain(T argc) {
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 @@ int main(int argc, char **argv) {
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
index 0272cd7c7577..3d686e4e8f5c 100644
--- 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 @@ T tmain(T argc) {
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 @@ int main(int argc, char **argv) {
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
index fe5330f2daec..2c8abf25d5ef 100644
--- a/clang/test/OpenMP/target_parallel_map_messages.cpp
+++ b/clang/test/OpenMP/target_parallel_map_messages.cpp
@@ -152,25 +152,25 @@ T tmain(T argc) {
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 @@ int main(int argc, char **argv) {
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
index 4a0522725800..ed0783ac90c9 100644
--- a/clang/test/OpenMP/target_simd_map_messages.cpp
+++ b/clang/test/OpenMP/target_simd_map_messages.cpp
@@ -146,25 +146,26 @@ T tmain(T argc) {
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 @@ int main(int argc, char **argv) {
#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
index 30dd2fded1b3..246228a4a98e 100644
--- a/clang/test/OpenMP/target_teams_distribute_map_messages.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_map_messages.cpp
@@ -152,25 +152,26 @@ T tmain(T argc) {
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 @@ int main(int argc, char **argv) {
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
index d82d9f65cb5b..082c2cb2b425 100644
--- 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 @@ T tmain(T argc) {
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 @@ int main(int argc, char **argv) {
#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
index 713103872016..90744b4cf0df 100644
--- 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 @@ T tmain(T argc) {
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 @@ int main(int argc, char **argv) {
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
index 0c47134494da..5229981b98fd 100644
--- 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 @@ T tmain(T argc) {
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 @@ int main(int argc, char **argv) {
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
index 95a79c591d7d..73e4cf59c5aa 100644
--- a/clang/test/OpenMP/target_teams_map_messages.cpp
+++ b/clang/test/OpenMP/target_teams_map_messages.cpp
@@ -62,7 +62,7 @@ struct SA {
#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 @@ void SAclient(int arg) {
{}
#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
diff erent 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
diff erent 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
diff erent 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
diff erent 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 @@ void SAclient(int arg) {
{}
#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 @@ void SAclient(int arg) {
#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 @@ T tmain(T argc) {
#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 @@ int main(int argc, char **argv) {
#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
index 1c9852f8f07a..911d2ae031eb 100644
--- a/clang/test/OpenMP/target_update_codegen.cpp
+++ b/clang/test/OpenMP/target_update_codegen.cpp
@@ -1113,7 +1113,7 @@ void array_shaping(float *f, int sa) {
// 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
index ddefdd884918..d2d0e0831d80 100644
--- a/openmp/libomptarget/src/omptarget.cpp
+++ b/openmp/libomptarget/src/omptarget.cpp
@@ -258,7 +258,7 @@ int targetDataMapper(ident_t *loc, DeviceTy &Device, void *arg_base, void *arg,
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
index 000000000000..87697affe06f
--- /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 <cstdio>
+#include <cstdlib>
+
+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<void *>(p) ? 1 : 0, spp[0][0].g[1],
+ spp[0][0].g == reinterpret_cast<void *>(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
+}
More information about the cfe-commits
mailing list