diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.h b/clang/lib/CodeGen/CGOpenMPRuntime.h
--- a/clang/lib/CodeGen/CGOpenMPRuntime.h
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.h
@@ -823,7 +823,7 @@
llvm::Value *Handle, llvm::Value *BasePtr,
llvm::Value *Ptr, llvm::Value *Size,
llvm::Value *MapType, CharUnits ElementSize,
- llvm::BasicBlock *ExitBB, bool IsInit);
+ llvm::BasicBlock *ExitBB);
struct TaskResultTy {
llvm::Value *NewTask = nullptr;
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7096,6 +7096,7 @@
std::pair HighestElem = {
0, Address::invalid()};
Address Base = Address::invalid();
+ bool HasCompleteRecord = false;
};
private:
@@ -7276,11 +7277,10 @@
if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_close)
!= MapModifiers.end())
Bits |= OMP_MAP_CLOSE;
- if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_present)
- != MapModifiers.end())
- Bits |= OMP_MAP_PRESENT;
- if (llvm::find(MotionModifiers, OMPC_MOTION_MODIFIER_present)
- != MotionModifiers.end())
+ if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_present) !=
+ MapModifiers.end() ||
+ llvm::find(MotionModifiers, OMPC_MOTION_MODIFIER_present) !=
+ MotionModifiers.end())
Bits |= OMP_MAP_PRESENT;
return Bits;
}
@@ -7816,6 +7816,10 @@
FirstPointerInComplexData = false;
}
}
+ // If ran into the whole component - allocate the space for the whole
+ // record.
+ if (!EncounteredME)
+ PartialStruct.HasCompleteRecord = true;
}
/// Return the adjusted map modifiers if the declaration a capture refers to
@@ -7924,137 +7928,97 @@
}
}
-public:
- MappableExprsHandler(const OMPExecutableDirective &Dir, CodeGenFunction &CGF)
- : CurDir(&Dir), CGF(CGF) {
- // Extract firstprivate clause information.
- for (const auto *C : Dir.getClausesOfKind())
- for (const auto *D : C->varlists())
- FirstPrivateDecls.try_emplace(
- cast(cast(D)->getDecl()), C->isImplicit());
- // Extract implicit firstprivates from uses_allocators clauses.
- for (const auto *C : Dir.getClausesOfKind()) {
- for (unsigned I = 0, E = C->getNumberOfAllocators(); I < E; ++I) {
- OMPUsesAllocatorsClause::Data D = C->getAllocatorData(I);
- if (const auto *DRE = dyn_cast_or_null(D.AllocatorTraits))
- FirstPrivateDecls.try_emplace(cast(DRE->getDecl()),
- /*Implicit=*/true);
- else if (const auto *VD = dyn_cast(
- cast(D.Allocator->IgnoreParenImpCasts())
- ->getDecl()))
- FirstPrivateDecls.try_emplace(VD, /*Implicit=*/true);
- }
- }
- // Extract device pointer clause information.
- for (const auto *C : Dir.getClausesOfKind())
- for (auto L : C->component_lists())
- DevPointersMap[std::get<0>(L)].push_back(std::get<1>(L));
- }
-
- /// Constructor for the declare mapper directive.
- MappableExprsHandler(const OMPDeclareMapperDecl &Dir, CodeGenFunction &CGF)
- : CurDir(&Dir), CGF(CGF) {}
-
- /// Generate code for the combined entry if we have a partially mapped struct
- /// and take care of the mapping flags of the arguments corresponding to
- /// individual struct members.
- void emitCombinedEntry(MapCombinedInfoTy &CombinedInfo,
- MapFlagsArrayTy &CurTypes,
- const StructRangeInfoTy &PartialStruct,
- bool NotTargetParams = false) const {
- // 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
- (*CurTypes.begin()) &= ~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, bool NotTargetParams = false,
+ void generateAllInfoForClauses(
+ ArrayRef Clauses, MapCombinedInfoTy &CombinedInfo,
const llvm::DenseSet> &SkipVarSet =
llvm::DenseSet>()) const {
// We have to process the component lists that relate with the same
// declaration in a single chunk so that we can generate the map flags
// correctly. Therefore, we organize all lists in a map.
- llvm::MapVector> Info;
+ enum MapKind { Present, Allocs, Other, Total };
+ llvm::MapVector,
+ SmallVector, 4>>
+ Info;
// Helper function to fill the information map for the different supported
// clauses.
auto &&InfoGen =
[&Info, &SkipVarSet](
- const ValueDecl *D,
+ const ValueDecl *D, MapKind Kind,
OMPClauseMappableExprCommon::MappableExprComponentListRef L,
OpenMPMapClauseKind MapType,
ArrayRef MapModifiers,
ArrayRef MotionModifiers,
bool ReturnDevicePointer, bool IsImplicit, const ValueDecl *Mapper,
bool ForDeviceAddr = false) {
- const ValueDecl *VD =
- D ? cast(D->getCanonicalDecl()) : nullptr;
- if (SkipVarSet.count(VD))
+ if (SkipVarSet.contains(D))
return;
- Info[VD].emplace_back(L, MapType, MapModifiers, MotionModifiers,
- ReturnDevicePointer, IsImplicit, Mapper,
- ForDeviceAddr);
+ auto It = Info.find(D);
+ if (It == Info.end())
+ It = Info
+ .insert(std::make_pair(
+ D, SmallVector, 4>(Total)))
+ .first;
+ It->second[Kind].emplace_back(L, MapType, MapModifiers,
+ MotionModifiers, ReturnDevicePointer,
+ IsImplicit, Mapper, ForDeviceAddr);
};
- assert(CurDir.is() &&
- "Expect a executable directive");
- const auto *CurExecDir = CurDir.get();
- for (const auto *C : CurExecDir->getClausesOfKind())
+ for (const auto *Cl : Clauses) {
+ const auto *C = dyn_cast(Cl);
+ if (!C)
+ continue;
+ MapKind Kind = Other;
+ if (!C->getMapTypeModifiers().empty() &&
+ llvm::any_of(C->getMapTypeModifiers(), [](OpenMPMapModifierKind K) {
+ return K == OMPC_MAP_MODIFIER_present;
+ }))
+ Kind = Present;
+ else if (C->getMapType() == OMPC_MAP_alloc)
+ Kind = Allocs;
for (const auto L : C->component_lists()) {
- 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));
}
- for (const auto *C : CurExecDir->getClausesOfKind())
+ }
+ for (const auto *Cl : Clauses) {
+ const auto *C = dyn_cast(Cl);
+ if (!C)
+ continue;
+ MapKind Kind = Other;
+ if (!C->getMotionModifiers().empty() &&
+ llvm::any_of(C->getMotionModifiers(), [](OpenMPMotionModifierKind K) {
+ return K == OMPC_MOTION_MODIFIER_present;
+ }))
+ Kind = Present;
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));
}
- for (const auto *C : CurExecDir->getClausesOfKind())
+ }
+ for (const auto *Cl : Clauses) {
+ const auto *C = dyn_cast(Cl);
+ if (!C)
+ continue;
+ MapKind Kind = Other;
+ if (!C->getMotionModifiers().empty() &&
+ llvm::any_of(C->getMotionModifiers(), [](OpenMPMotionModifierKind K) {
+ return K == OMPC_MOTION_MODIFIER_present;
+ }))
+ Kind = Present;
for (const auto L : C->component_lists()) {
- InfoGen(std::get<0>(L), std::get<1>(L), OMPC_MAP_from, llvm::None,
- C->getMotionModifiers(), /*ReturnDevicePointer=*/false,
- C->isImplicit(), std::get<2>(L));
+ 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));
}
+ }
// Look at the use_device_ptr clause information and mark the existing map
// entries as such. If there is no map information for an entry in the
@@ -8062,12 +8026,15 @@
// section. It is the user fault if that was not mapped before. If there is
// no map information and the pointer is a struct member, then we defer the
// emission of that entry until the whole struct has been processed.
- llvm::MapVector>
+ llvm::MapVector,
+ SmallVector>
DeferredInfo;
MapCombinedInfoTy UseDevicePtrCombinedInfo;
- for (const auto *C :
- CurExecDir->getClausesOfKind()) {
+ for (const auto *Cl : Clauses) {
+ const auto *C = dyn_cast(Cl);
+ if (!C)
+ continue;
for (const auto L : C->component_lists()) {
OMPClauseMappableExprCommon::MappableExprComponentListRef Components =
std::get<1>(L);
@@ -8084,28 +8051,34 @@
// We potentially have map information for this declaration already.
// Look for the first set of components that refer to it.
if (It != Info.end()) {
- auto *CI = llvm::find_if(It->second, [VD](const MapInfo &MI) {
- return MI.Components.back().getAssociatedDeclaration() == VD;
- });
- // If we found a map entry, signal that the pointer has to be returned
- // and move on to the next declaration.
- // Exclude cases where the base pointer is mapped as array subscript,
- // array section or array shaping. The base address is passed as a
- // pointer to base in this case and cannot be used as a base for
- // use_device_ptr list item.
- if (CI != It->second.end()) {
- auto PrevCI = std::next(CI->Components.rbegin());
- const auto *VarD = dyn_cast(VD);
- if (CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory() ||
- isa(IE) ||
- !VD->getType().getNonReferenceType()->isPointerType() ||
- PrevCI == CI->Components.rend() ||
- isa(PrevCI->getAssociatedExpression()) || !VarD ||
- VarD->hasLocalStorage()) {
- CI->ReturnDevicePointer = true;
- continue;
+ bool Found = false;
+ for (auto &Data : It->second) {
+ auto *CI = llvm::find_if(Data, [VD](const MapInfo &MI) {
+ return MI.Components.back().getAssociatedDeclaration() == VD;
+ });
+ // If we found a map entry, signal that the pointer has to be
+ // returned and move on to the next declaration. Exclude cases where
+ // the base pointer is mapped as array subscript, array section or
+ // array shaping. The base address is passed as a pointer to base in
+ // this case and cannot be used as a base for use_device_ptr list
+ // item.
+ if (CI != Data.end()) {
+ auto PrevCI = std::next(CI->Components.rbegin());
+ const auto *VarD = dyn_cast(VD);
+ if (CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory() ||
+ isa(IE) ||
+ !VD->getType().getNonReferenceType()->isPointerType() ||
+ PrevCI == CI->Components.rend() ||
+ isa(PrevCI->getAssociatedExpression()) || !VarD ||
+ VarD->hasLocalStorage()) {
+ CI->ReturnDevicePointer = true;
+ Found = true;
+ break;
+ }
}
}
+ if (Found)
+ continue;
}
// We didn't find any match in our map information - generate a zero
@@ -8119,8 +8092,9 @@
// Nonetheless, generateInfoForComponentList must be called to take
// the pointer into account for the calculation of the range of the
// partial struct.
- InfoGen(nullptr, Components, OMPC_MAP_unknown, llvm::None, llvm::None,
- /*ReturnDevicePointer=*/false, C->isImplicit(), nullptr);
+ InfoGen(nullptr, Other, Components, OMPC_MAP_unknown, llvm::None,
+ llvm::None, /*ReturnDevicePointer=*/false, C->isImplicit(),
+ nullptr);
DeferredInfo[nullptr].emplace_back(IE, VD, /*ForDeviceAddr=*/false);
} else {
llvm::Value *Ptr =
@@ -8129,9 +8103,7 @@
UseDevicePtrCombinedInfo.Pointers.push_back(Ptr);
UseDevicePtrCombinedInfo.Sizes.push_back(
llvm::Constant::getNullValue(CGF.Int64Ty));
- UseDevicePtrCombinedInfo.Types.push_back(
- OMP_MAP_RETURN_PARAM |
- (NotTargetParams ? OMP_MAP_NONE : OMP_MAP_TARGET_PARAM));
+ UseDevicePtrCombinedInfo.Types.push_back(OMP_MAP_RETURN_PARAM);
UseDevicePtrCombinedInfo.Mappers.push_back(nullptr);
}
}
@@ -8144,8 +8116,10 @@
// no map information and the pointer is a struct member, then we defer the
// emission of that entry until the whole struct has been processed.
llvm::SmallDenseSet, 4> Processed;
- for (const auto *C :
- CurExecDir->getClausesOfKind()) {
+ for (const auto *Cl : Clauses) {
+ const auto *C = dyn_cast(Cl);
+ if (!C)
+ continue;
for (const auto L : C->component_lists()) {
assert(!std::get<1>(L).empty() &&
"Not expecting empty list of components!");
@@ -8162,15 +8136,21 @@
// We potentially have map information for this declaration already.
// Look for the first set of components that refer to it.
if (It != Info.end()) {
- auto *CI = llvm::find_if(It->second, [VD](const MapInfo &MI) {
- return MI.Components.back().getAssociatedDeclaration() == VD;
- });
- // If we found a map entry, signal that the pointer has to be returned
- // and move on to the next declaration.
- if (CI != It->second.end()) {
- CI->ReturnDevicePointer = true;
- continue;
+ bool Found = false;
+ for (auto &Data : It->second) {
+ auto *CI = llvm::find_if(Data, [VD](const MapInfo &MI) {
+ return MI.Components.back().getAssociatedDeclaration() == VD;
+ });
+ // If we found a map entry, signal that the pointer has to be
+ // returned and move on to the next declaration.
+ if (CI != Data.end()) {
+ CI->ReturnDevicePointer = true;
+ Found = true;
+ break;
+ }
}
+ if (Found)
+ continue;
}
// We didn't find any match in our map information - generate a zero
@@ -8184,7 +8164,7 @@
// Nonetheless, generateInfoForComponentList must be called to take
// the pointer into account for the calculation of the range of the
// partial struct.
- InfoGen(nullptr, std::get<1>(L), OMPC_MAP_unknown, llvm::None,
+ InfoGen(nullptr, Other, std::get<1>(L), OMPC_MAP_unknown, llvm::None,
llvm::None, /*ReturnDevicePointer=*/false, C->isImplicit(),
nullptr, /*ForDeviceAddr=*/true);
DeferredInfo[nullptr].emplace_back(IE, VD, /*ForDeviceAddr=*/true);
@@ -8198,55 +8178,49 @@
CombinedInfo.Pointers.push_back(Ptr);
CombinedInfo.Sizes.push_back(
llvm::Constant::getNullValue(CGF.Int64Ty));
- CombinedInfo.Types.push_back(
- OMP_MAP_RETURN_PARAM |
- (NotTargetParams ? OMP_MAP_NONE : OMP_MAP_TARGET_PARAM));
+ CombinedInfo.Types.push_back(OMP_MAP_RETURN_PARAM);
CombinedInfo.Mappers.push_back(nullptr);
}
}
}
- 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 = !NotTargetParams;
-
+ 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();
- generateInfoForComponentList(L.MapType, L.MapModifiers,
- L.MotionModifiers, L.Components, CurInfo,
- PartialStruct, IsFirstComponentList,
- L.IsImplicit, L.Mapper, L.ForDeviceAddr);
-
- // 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;
+ 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();
+ generateInfoForComponentList(
+ L.MapType, L.MapModifiers, L.MotionModifiers, L.Components,
+ CurInfo, PartialStruct, /*IsFirstComponentList=*/false,
+ L.IsImplicit, L.Mapper, L.ForDeviceAddr);
+
+ // 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;
+ }
}
- IsFirstComponentList = false;
}
// 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;
@@ -8265,9 +8239,9 @@
BasePtr = this->CGF.EmitLValue(L.IE).getPointer(CGF);
Ptr = this->CGF.EmitLoadOfScalar(this->CGF.EmitLValue(L.IE),
L.IE->getExprLoc());
- // Entry is PTR_AND_OBJ and RETURN_PARAM. Also, set the placeholder
- // value MEMBER_OF=FFFF so that the entry is later updated with the
- // correct value of MEMBER_OF.
+ // Entry is PTR_AND_OBJ and RETURN_PARAM. Also, set the
+ // placeholder value MEMBER_OF=FFFF so that the entry is later
+ // updated with the correct value of MEMBER_OF.
CurInfo.Types.push_back(OMP_MAP_PTR_AND_OBJ | OMP_MAP_RETURN_PARAM |
OMP_MAP_MEMBER_OF);
}
@@ -8278,74 +8252,124 @@
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())
- emitCombinedEntry(CombinedInfo, CurInfo.Types, PartialStruct,
- NotTargetParams);
+ emitCombinedEntry(CombinedInfo, CurInfo.Types, PartialStruct);
- // We need to append the results of this capture to what we already have.
+ // We need to append the results of this capture to what we already
+ // have.
CombinedInfo.append(CurInfo);
}
// Append data for use_device_ptr clauses.
CombinedInfo.append(UseDevicePtrCombinedInfo);
}
- /// Generate all the base pointers, section pointers, sizes, map types, and
- /// mappers for the extracted map clauses of user-defined mapper (all included
- /// in \a CombinedInfo).
- void generateAllInfoForMapper(MapCombinedInfoTy &CombinedInfo) const {
- assert(CurDir.is() &&
- "Expect a declare mapper directive");
- const auto *CurMapperDir = CurDir.get();
- // We have to process the component lists that relate with the same
- // declaration in a single chunk so that we can generate the map flags
- // correctly. Therefore, we organize all lists in a map.
- llvm::MapVector> Info;
-
- // Fill the information map for map clauses.
- for (const auto *C : CurMapperDir->clauselists()) {
- const auto *MC = cast(C);
- for (const auto L : MC->component_lists()) {
- const ValueDecl *VD =
- std::get<0>(L) ? cast(std::get<0>(L)->getCanonicalDecl())
- : nullptr;
- // Get the corresponding user-defined mapper.
- Info[VD].emplace_back(std::get<1>(L), MC->getMapType(),
- MC->getMapTypeModifiers(), llvm::None,
- /*ReturnDevicePointer=*/false, MC->isImplicit(),
- std::get<2>(L));
+public:
+ MappableExprsHandler(const OMPExecutableDirective &Dir, CodeGenFunction &CGF)
+ : CurDir(&Dir), CGF(CGF) {
+ // Extract firstprivate clause information.
+ for (const auto *C : Dir.getClausesOfKind())
+ for (const auto *D : C->varlists())
+ FirstPrivateDecls.try_emplace(
+ cast(cast(D)->getDecl()), C->isImplicit());
+ // Extract implicit firstprivates from uses_allocators clauses.
+ for (const auto *C : Dir.getClausesOfKind()) {
+ for (unsigned I = 0, E = C->getNumberOfAllocators(); I < E; ++I) {
+ OMPUsesAllocatorsClause::Data D = C->getAllocatorData(I);
+ if (const auto *DRE = dyn_cast_or_null(D.AllocatorTraits))
+ FirstPrivateDecls.try_emplace(cast(DRE->getDecl()),
+ /*Implicit=*/true);
+ else if (const auto *VD = dyn_cast(
+ cast(D.Allocator->IgnoreParenImpCasts())
+ ->getDecl()))
+ FirstPrivateDecls.try_emplace(VD, /*Implicit=*/true);
}
}
+ // Extract device pointer clause information.
+ for (const auto *C : Dir.getClausesOfKind())
+ for (auto L : C->component_lists())
+ DevPointersMap[std::get<0>(L)].push_back(std::get<1>(L));
+ }
- for (const auto &M : Info) {
- // We need to know when we generate information for the first component
- // associated with a capture, because the mapping flags depend on it.
- bool IsFirstComponentList = true;
+ /// 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,
+ bool NotTargetParams = true) const {
+ Address LBAddr = PartialStruct.LowestElem.second;
+ Address HBAddr = PartialStruct.HighestElem.second;
+ if (PartialStruct.HasCompleteRecord) {
+ LBAddr = PartialStruct.Base;
+ HBAddr = PartialStruct.Base;
+ }
+ // 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);
- 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())
- emitCombinedEntry(CombinedInfo, CurInfo.Types, PartialStruct);
+ /// Generate all the base pointers, section pointers, sizes, map types, and
+ /// mappers for the extracted mappable expressions (all included in \a
+ /// CombinedInfo). Also, for each item that relates with a device pointer, a
+ /// pair of the relevant declaration and index where it occurs is appended to
+ /// the device pointers info array.
+ void generateAllInfo(
+ MapCombinedInfoTy &CombinedInfo,
+ const llvm::DenseSet> &SkipVarSet =
+ llvm::DenseSet>()) const {
+ assert(CurDir.is() &&
+ "Expect a executable directive");
+ const auto *CurExecDir = CurDir.get();
+ generateAllInfoForClauses(CurExecDir->clauses(), CombinedInfo, SkipVarSet);
+ }
- // We need to append the results of this capture to what we already have.
- CombinedInfo.append(CurInfo);
- }
+ /// Generate all the base pointers, section pointers, sizes, map types, and
+ /// mappers for the extracted map clauses of user-defined mapper (all included
+ /// in \a CombinedInfo).
+ void generateAllInfoForMapper(MapCombinedInfoTy &CombinedInfo) const {
+ assert(CurDir.is() &&
+ "Expect a declare mapper directive");
+ const auto *CurMapperDir = CurDir.get();
+ generateAllInfoForClauses(CurMapperDir->clauses(), CombinedInfo);
}
/// Emit capture info for lambdas for variables captured by reference.
@@ -8491,6 +8515,25 @@
C->isImplicit(), Mapper);
}
}
+ llvm::stable_sort(DeclComponentLists, [](const MapData &LHS,
+ const MapData &RHS) {
+ ArrayRef MapModifiers = std::get<2>(LHS);
+ OpenMPMapClauseKind MapType = std::get<1>(RHS);
+ bool HasPresent = !MapModifiers.empty() &&
+ llvm::any_of(MapModifiers, [](OpenMPMapModifierKind K) {
+ return K == clang::OMPC_MAP_MODIFIER_present;
+ });
+ bool HasAllocs = MapType == OMPC_MAP_alloc;
+ MapModifiers = std::get<2>(RHS);
+ MapType = std::get<1>(LHS);
+ bool HasPresentR =
+ !MapModifiers.empty() &&
+ llvm::any_of(MapModifiers, [](OpenMPMapModifierKind K) {
+ return K == clang::OMPC_MAP_MODIFIER_present;
+ });
+ bool HasAllocsR = MapType == OMPC_MAP_alloc;
+ return (HasPresent && !HasPresentR) || (HasAllocs && !HasAllocsR);
+ });
// Find overlapping elements (including the offset from the base element).
llvm::SmallDenseMap<
@@ -8548,7 +8591,7 @@
}
}
for (auto &Pair : OverlappedData) {
- llvm::sort(
+ llvm::stable_sort(
Pair.getSecond(),
[&Layout](
OMPClauseMappableExprCommon::MappableExprComponentListRef First,
@@ -8590,6 +8633,7 @@
// Associated with a capture, because the mapping flags depend on it.
// Go through all of the elements with the overlapped elements.
+ bool IsFirstComponentList = true;
for (const auto &Pair : OverlappedData) {
const MapData &L = *Pair.getFirst();
OMPClauseMappableExprCommon::MappableExprComponentListRef Components;
@@ -8600,14 +8644,13 @@
std::tie(Components, MapType, MapModifiers, IsImplicit, Mapper) = L;
ArrayRef
OverlappedComponents = Pair.getSecond();
- bool IsFirstComponentList = true;
generateInfoForComponentList(
MapType, MapModifiers, llvm::None, Components, CombinedInfo,
PartialStruct, IsFirstComponentList, IsImplicit, Mapper,
/*ForDeviceAddr=*/false, OverlappedComponents);
+ IsFirstComponentList = false;
}
// Go through other elements without overlapped elements.
- bool IsFirstComponentList = OverlappedData.empty();
for (const MapData &L : DeclComponentLists) {
OMPClauseMappableExprCommon::MappableExprComponentListRef Components;
OpenMPMapClauseKind MapType;
@@ -9039,10 +9082,6 @@
/// void .omp_mapper...(void *rt_mapper_handle,
/// void *base, void *begin,
/// int64_t size, int64_t type) {
-/// // Allocate space for an array section first.
-/// if (size > 1 && !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:
@@ -9053,12 +9092,12 @@
/// else
/// __tgt_push_mapper_component(rt_mapper_handle, c.arg_base,
/// c.arg_begin, c.arg_size, c.arg_type);
-/// }
-/// }
-/// // Delete the array section.
-/// if (size > 1 && maptype.IsDelete)
+/// // Delete/allocate space for an array section first.
+/// if (size > 1)
/// __tgt_push_mapper_component(rt_mapper_handle, base, begin,
/// size*sizeof(Ty), clearToFrom(type));
+/// }
+/// }
/// }
/// \endcode
void CGOpenMPRuntime::emitUserDefinedMapper(const OMPDeclareMapperDecl *D,
@@ -9130,16 +9169,6 @@
MapperCGF.GetAddrOfLocalVar(&BeginArg),
/*Volatile=*/false, C.getPointerType(C.VoidPtrTy), Loc);
- // Emit array initiation if this is an array section and \p MapType indicates
- // that memory allocation is required.
- llvm::BasicBlock *HeadBB = MapperCGF.createBasicBlock("omp.arraymap.head");
- emitUDMapperArrayInitOrDel(MapperCGF, Handle, BaseIn, BeginIn, Size, MapType,
- ElementSize, HeadBB, /*IsInit=*/true);
-
- // Emit a for loop to iterate through SizeArg of elements and map all of them.
-
- // Emit the loop header block.
- MapperCGF.EmitBlock(HeadBB);
llvm::BasicBlock *BodyBB = MapperCGF.createBasicBlock("omp.arraymap.body");
llvm::BasicBlock *DoneBB = MapperCGF.createBasicBlock("omp.done");
// Evaluate whether the initial condition is satisfied.
@@ -9310,7 +9339,7 @@
// Emit array deletion if this is an array section and \p MapType indicates
// that deletion is required.
emitUDMapperArrayInitOrDel(MapperCGF, Handle, BaseIn, BeginIn, Size, MapType,
- ElementSize, DoneBB, /*IsInit=*/false);
+ ElementSize, DoneBB);
// Emit the function exit block.
MapperCGF.EmitBlock(DoneBB, /*IsFinished=*/true);
@@ -9331,32 +9360,12 @@
void CGOpenMPRuntime::emitUDMapperArrayInitOrDel(
CodeGenFunction &MapperCGF, llvm::Value *Handle, llvm::Value *Base,
llvm::Value *Begin, llvm::Value *Size, llvm::Value *MapType,
- CharUnits ElementSize, llvm::BasicBlock *ExitBB, bool IsInit) {
- StringRef Prefix = IsInit ? ".init" : ".del";
-
+ CharUnits ElementSize, llvm::BasicBlock *ExitBB) {
// 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::BasicBlock *BodyBB = MapperCGF.createBasicBlock("omp.array.body");
llvm::Value *IsArray = MapperCGF.Builder.CreateICmpSGE(
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);
- llvm::Value *DeleteBit = MapperCGF.Builder.CreateAnd(
- MapType,
- MapperCGF.Builder.getInt64(MappableExprsHandler::OMP_MAP_DELETE));
- llvm::Value *DeleteCond;
- if (IsInit) {
- DeleteCond = MapperCGF.Builder.CreateIsNull(
- DeleteBit, getName({"omp.array", Prefix, ".delete"}));
- } else {
- DeleteCond = MapperCGF.Builder.CreateIsNotNull(
- DeleteBit, getName({"omp.array", Prefix, ".delete"}));
- }
- MapperCGF.Builder.CreateCondBr(DeleteCond, BodyBB, ExitBB);
+ MapperCGF.Builder.CreateCondBr(IsArray, BodyBB, ExitBB);
MapperCGF.EmitBlock(BodyBB);
// Get the array size by multiplying element size and element number (i.e., \p
@@ -9599,7 +9608,7 @@
llvm::DenseSet> MappedVarSet;
auto RI = CS.getCapturedRecordDecl()->field_begin();
- auto CV = CapturedVars.begin();
+ auto *CV = CapturedVars.begin();
for (CapturedStmt::const_capture_iterator CI = CS.capture_begin(),
CE = CS.capture_end();
CI != CE; ++CI, ++RI, ++CV) {
@@ -9646,7 +9655,8 @@
// 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())
- MEHandler.emitCombinedEntry(CombinedInfo, CurInfo.Types, PartialStruct);
+ MEHandler.emitCombinedEntry(CombinedInfo, CurInfo.Types, PartialStruct,
+ /*NoTargetParam=*/false);
// We need to append the results of this capture to what we already have.
CombinedInfo.append(CurInfo);
@@ -9657,8 +9667,7 @@
CombinedInfo.Types);
// Map any list items in a map clause that were not captures because they
// weren't referenced within the construct.
- MEHandler.generateAllInfo(CombinedInfo, /*NotTargetParams=*/true,
- MappedVarSet);
+ MEHandler.generateAllInfo(CombinedInfo, MappedVarSet);
TargetDataInfo Info;
// Fill up the arrays and create the arguments.
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -3471,9 +3471,11 @@
!Stack->isLoopControlVariable(VD).first) {
if (!Stack->checkMappableExprComponentListsForDecl(
VD, /*CurrentRegionOnly=*/true,
- [](OMPClauseMappableExprCommon::MappableExprComponentListRef
- StackComponents,
- OpenMPClauseKind) {
+ [this](OMPClauseMappableExprCommon::MappableExprComponentListRef
+ StackComponents,
+ OpenMPClauseKind) {
+ if (SemaRef.LangOpts.OpenMP >= 50)
+ return !StackComponents.empty();
// Variable is used if it has been marked as an array, array
// section, array shaping or the variable iself.
return StackComponents.size() == 1 ||
@@ -17123,7 +17125,9 @@
ERange, CKind, &EnclosingExpr,
CurComponents](OMPClauseMappableExprCommon::MappableExprComponentListRef
StackComponents,
- OpenMPClauseKind) {
+ OpenMPClauseKind Kind) {
+ if (CKind == Kind && SemaRef.LangOpts.OpenMP >= 50)
+ return false;
assert(!StackComponents.empty() &&
"Map clause expression with no components!");
assert(StackComponents.back().getAssociatedDeclaration() == VD &&
diff --git a/clang/test/OpenMP/declare_mapper_codegen.cpp b/clang/test/OpenMP/declare_mapper_codegen.cpp
--- a/clang/test/OpenMP/declare_mapper_codegen.cpp
+++ b/clang/test/OpenMP/declare_mapper_codegen.cpp
@@ -37,25 +37,25 @@
// CK0: [[TEAMNWTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 33]
// CK0-64: [[EDSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
// CK0-32: [[EDSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
-// CK0: [[EDTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK0: [[EDTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 1]
// CK0-64: [[EDNWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
// CK0-32: [[EDNWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
-// CK0: [[EDNWTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK0: [[EDNWTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 1]
// CK0-64: [[EXDSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
// CK0-32: [[EXDSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
-// CK0: [[EXDTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 34]
+// CK0: [[EXDTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 2]
// CK0-64: [[EXDNWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
// CK0-32: [[EXDNWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
-// CK0: [[EXDNWTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 34]
+// CK0: [[EXDNWTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 2]
// CK0-64: [[TSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
// CK0-32: [[TSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
-// CK0: [[TTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK0: [[TTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 1]
// CK0-64: [[FSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
// CK0-32: [[FSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
-// CK0: [[FTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 34]
+// CK0: [[FTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 2]
// CK0-64: [[FNWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
// CK0-32: [[FNWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
-// CK0: [[FNWTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 34]
+// CK0: [[FNWTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 2]
class C {
public:
@@ -80,25 +80,10 @@
// 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: [[INIT]]
-// CK0-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
-// CK0-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
-// CK0-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
-// CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]])
-// CK0: br label %[[LHEAD:[^,]+]]
-
-// CK0: [[LHEAD]]
// 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: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %{{.+}} ], [ [[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
@@ -120,10 +105,10 @@
// CK0-DAG: [[MEMBER]]
// CK0-DAG: br i1 true, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
// CK0-DAG: [[MEMBERCOM]]
-// CK0-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 32, [[SHIPRESIZE]]
+// CK0-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 0, [[SHIPRESIZE]]
// CK0-DAG: br label %[[LTYPE]]
// CK0-DAG: [[LTYPE]]
-// CK0-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 32, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
+// CK0-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 0, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
// CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
// CK0-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
// CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
@@ -217,10 +202,6 @@
// CK0: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1
// CK0: br i1 [[ISARRAY]], label %[[EVALDEL:[^,]+]], label %[[DONE]]
// CK0: [[EVALDEL]]
-// CK0: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
-// CK0: [[ISDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
-// CK0: br i1 [[ISDEL]], label %[[DEL:[^,]+]], label %[[DONE]]
-// CK0: [[DEL]]
// 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
@@ -460,24 +441,10 @@
// 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: [[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: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]]
// CK1: [[LBODY]]
-// CK1: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %[[LHEAD]] ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
+// CK1: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %{{.+}} ], [ [[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-DAG: [[AEND:%.+]] = getelementptr i32, i32* [[ABEGIN]], i32 1
@@ -495,10 +462,10 @@
// CK1-DAG: [[MEMBER]]
// CK1-DAG: br i1 true, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
// CK1-DAG: [[MEMBERCOM]]
-// CK1-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 32, [[SHIPRESIZE]]
+// CK1-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 0, [[SHIPRESIZE]]
// CK1-DAG: br label %[[LTYPE]]
// CK1-DAG: [[LTYPE]]
-// CK1-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 32, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
+// CK1-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 0, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
// CK1-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
// CK1-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
// CK1-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
@@ -560,10 +527,6 @@
// CK1: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1
// CK1: br i1 [[ISARRAY]], label %[[EVALDEL:[^,]+]], label %[[DONE]]
// CK1: [[EVALDEL]]
-// CK1: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
-// CK1: [[ISDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
-// CK1: br i1 [[ISDEL]], label %[[DEL:[^,]+]], label %[[DONE]]
-// CK1: [[DEL]]
// 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]])
@@ -623,24 +586,10 @@
// 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: [[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: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]]
// CK2: [[LBODY]]
-// CK2: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %[[LHEAD]] ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
+// CK2: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %{{.+}} ], [ [[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-DAG: [[BEND:%.+]] = getelementptr %class.B, %class.B* [[BBEGIN]], i32 1
@@ -658,10 +607,10 @@
// CK2-DAG: [[MEMBER]]
// CK2-DAG: br i1 true, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
// CK2-DAG: [[MEMBERCOM]]
-// CK2-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 32, [[SHIPRESIZE]]
+// CK2-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 0, [[SHIPRESIZE]]
// CK2-DAG: br label %[[LTYPE]]
// CK2-DAG: [[LTYPE]]
-// CK2-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 32, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
+// CK2-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 0, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
// CK2-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
// CK2-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
// CK2-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
@@ -723,10 +672,6 @@
// CK2: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1
// CK2: br i1 [[ISARRAY]], label %[[EVALDEL:[^,]+]], label %[[DONE]]
// CK2: [[EVALDEL]]
-// CK2: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
-// CK2: [[ISDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
-// CK2: br i1 [[ISDEL]], label %[[DEL:[^,]+]], label %[[DONE]]
-// CK2: [[DEL]]
// 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]])
@@ -859,13 +804,13 @@
// CK4-64: [[TSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
// CK4-32: [[TSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
-// PRESENT=0x1000 | TARGET_PARAM=0x20 | TO=0x1 = 0x1021
-// CK4: [[TTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1021]]]
+// PRESENT=0x1000 | TO=0x1 = 0x1001
+// CK4: [[TTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1001]]]
// CK4-64: [[FSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
// CK4-32: [[FSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
-// PRESENT=0x1000 | TARGET_PARAM=0x20 | FROM=0x2 = 0x1022
-// CK4: [[FTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1022]]]
+// PRESENT=0x1000 | FROM=0x2 = 0x1002
+// CK4: [[FTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1002]]]
class C {
public:
@@ -890,25 +835,10 @@
// 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: [[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
-// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]])
-// CK4: br label %[[LHEAD:[^,]+]]
-
-// CK4: [[LHEAD]]
// 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: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %{{.+}} ], [ [[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
@@ -930,10 +860,10 @@
// CK4-DAG: [[MEMBER]]
// CK4-DAG: br i1 true, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
// CK4-DAG: [[MEMBERCOM]]
-// CK4-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 32, [[SHIPRESIZE]]
+// CK4-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 0, [[SHIPRESIZE]]
// CK4-DAG: br label %[[LTYPE]]
// CK4-DAG: [[LTYPE]]
-// CK4-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 32, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
+// CK4-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 0, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
// CK4-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
// CK4-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
// CK4-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
@@ -1027,10 +957,6 @@
// CK4: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1
// CK4: br i1 [[ISARRAY]], label %[[EVALDEL:[^,]+]], label %[[DONE]]
// CK4: [[EVALDEL]]
-// CK4: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
-// CK4: [[ISDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
-// CK4: br i1 [[ISDEL]], label %[[DEL:[^,]+]], label %[[DONE]]
-// CK4: [[DEL]]
// CK4-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
// CK4-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
// CK4-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
diff --git a/clang/test/OpenMP/target_data_codegen.cpp b/clang/test/OpenMP/target_data_codegen.cpp
--- a/clang/test/OpenMP/target_data_codegen.cpp
+++ b/clang/test/OpenMP/target_data_codegen.cpp
@@ -30,19 +30,19 @@
double gc[100];
// CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 800]
-// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 34]
+// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 2]
// CK1: [[SIZE02:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 4]
-// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 1]
-// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 37]
+// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 5]
// CK1: [[SIZE04:@.+]] = {{.+}}constant [2 x i64] [i64 sdiv exact (i64 sub (i64 ptrtoint (double** getelementptr (double*, double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1), i32 1) to i64), i64 ptrtoint (double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1) to i64)), i64 ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)), i64 24]
-// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673]
+// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710673]
-// CK1: [[MTYPE05:@.+]] = {{.+}}constant [1 x i64] [i64 1057]
+// CK1: [[MTYPE05:@.+]] = {{.+}}constant [1 x i64] [i64 1025]
-// CK1: [[MTYPE06:@.+]] = {{.+}}constant [1 x i64] [i64 1061]
+// CK1: [[MTYPE06:@.+]] = {{.+}}constant [1 x i64] [i64 1029]
// CK1-LABEL: _Z3fooi
void foo(int arg) {
@@ -255,17 +255,17 @@
ST gb;
double gc[100];
-// PRESENT=0x1000 | TARGET_PARAM=0x20 | TO=0x1 = 0x1021
-// CK1A: [[MTYPE00Begin:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1021]]]
+// PRESENT=0x1000 | TO=0x1 = 0x1001
+// CK1A: [[MTYPE00Begin:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1001]]]
-// TARGET_PARAM=0x20 | TO=0x1 = 0x21
-// CK1A: [[MTYPE00End:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x21]]]
+// TO=0x1 = 0x1
+// CK1A: [[MTYPE00End:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1]]]
-// PRESENT=0x1000 | CLOSE=0x400 | TARGET_PARAM=0x20 | ALWAYS=0x4 | TO=0x1 = 0x1425
-// CK1A: [[MTYPE01Begin:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1425]]]
+// PRESENT=0x1000 | CLOSE=0x400 | ALWAYS=0x4 | TO=0x1 = 0x14205
+// CK1A: [[MTYPE01Begin:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1405]]]
-// CLOSE=0x400 | TARGET_PARAM=0x20 | ALWAYS=0x4 | TO=0x1 = 0x425
-// CK1A: [[MTYPE01End:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x425]]]
+// CLOSE=0x400 | ALWAYS=0x4 | TO=0x1 = 0x405
+// CK1A: [[MTYPE01End:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x405]]]
// CK1A-LABEL: _Z3fooi
void foo(int arg) {
@@ -357,7 +357,7 @@
}
};
-// CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710677]
+// CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710677]
// CK2-LABEL: _Z3bari
int bar(int arg){
@@ -475,7 +475,7 @@
}
};
-// CK4: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976711701]
+// CK4: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976711701]
// CK4-LABEL: _Z3bari
int bar(int arg){
@@ -561,8 +561,9 @@
void test_close_modifier(int arg) {
S2 *ps;
- // CK5: private unnamed_addr constant [5 x i64] [i64 1059, i64 32, i64 562949953421328, i64 16, i64 1043]
- #pragma omp target data map(close,tofrom: arg, ps->ps->ps->ps->s)
+// CK5: private unnamed_addr constant [5 x i64] [i64 1027, i64 0, i64 562949953421328, i64 16, i64 1043]
+#pragma omp target data map(close, tofrom \
+ : arg, ps->ps->ps->ps->s)
{
++(arg);
}
@@ -585,8 +586,9 @@
// SIMD-ONLY2-NOT: {{__kmpc|__tgt}}
#ifdef CK6
void test_close_modifier(int arg) {
- // CK6: private unnamed_addr constant [1 x i64] [i64 1059]
- #pragma omp target data map(close,tofrom: arg)
+// CK6: private unnamed_addr constant [1 x i64] [i64 1027]
+#pragma omp target data map(close, tofrom \
+ : arg)
{++arg;}
}
#endif
@@ -642,36 +644,39 @@
// CK8: private unnamed_addr constant [11 x i64]
- // ps1
- //
- // PRESENT=0x1000 | TARGET_PARAM=0x20 = 0x1020
- // 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 [[#0x1020]], i64 [[#0x1000000000003]],
- // CK8-SAME: {{^}} i64 [[#0x1000000001010]], i64 [[#0x1010]], i64 [[#0x1013]],
-
- // arg
- //
- // PRESENT=0x1000 | TARGET_PARAM=0x20 | FROM=0x2 | TO=0x1 = 0x1023
- //
- // CK8-SAME: {{^}} i64 [[#0x1023]],
-
- // ps2
- //
- // PRESENT=0x1000 | TARGET_PARAM=0x20 = 0x1020
- // 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 [[#0x1020]], i64 [[#0x7000000001003]],
- // CK8-SAME: {{^}} i64 [[#0x7000000000010]], i64 [[#0x10]], i64 [[#0x13]]]
- #pragma omp target data map(tofrom: ps1->s) \
- map(present,tofrom: arg, ps1->ps->ps->ps->s, ps2->s) \
- map(tofrom: ps2->ps->ps->ps->s)
+// ps1
+//
+// PRESENT=0x1000 = 0x1000
+// MEMBER_OF_1=0x1000000000000 | PRESENT=0x1000 | PTR_AND_OBJ=0x10 = 0x1000000001010
+// PRESENT=0x1000 | PTR_AND_OBJ=0x10 = 0x1010
+// PRESENT=0x1000 | PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x1013
+// MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003
+//
+// CK8-SAME: {{^}} [i64 [[#0x1000]], i64 [[#0x1000000001010]],
+// CK8-SAME: {{^}} i64 [[#0x1010]], i64 [[#0x1013]], i64 [[#0x1000000000003]],
+
+// arg
+//
+// PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x1003
+//
+// CK8-SAME: {{^}} i64 [[#0x1003]],
+
+// ps2
+//
+// PRESENT=0x1000 = 0x1000
+// MEMBER_OF_7=0x7000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x7000000001003
+// MEMBER_OF_7=0x7000000000000 | PTR_AND_OBJ=0x10 = 0x7000000000010
+// PTR_AND_OBJ=0x10 = 0x10
+// PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x13
+//
+// CK8-SAME: {{^}} i64 [[#0x1000]], i64 [[#0x7000000001003]],
+// CK8-SAME: {{^}} i64 [[#0x7000000000010]], i64 [[#0x10]], i64 [[#0x13]]]
+#pragma omp target data map(tofrom \
+ : ps1->s) \
+ map(present, tofrom \
+ : arg, ps1->ps->ps->ps->s, ps2->s) \
+ map(tofrom \
+ : ps2->ps->ps->ps->s)
{
++(arg);
}
@@ -694,9 +699,10 @@
// SIMD-ONLY2-NOT: {{__kmpc|__tgt}}
#ifdef CK9
void test_present_modifier(int arg) {
- // PRESENT=0x1000 | TARGET_PARAM=0x20 | FROM=0x2 | TO=0x1 = 0x1023
- // CK9: private unnamed_addr constant [1 x i64] [i64 [[#0x1023]]]
- #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_data_use_device_addr_codegen.cpp b/clang/test/OpenMP/target_data_use_device_addr_codegen.cpp
--- a/clang/test/OpenMP/target_data_use_device_addr_codegen.cpp
+++ b/clang/test/OpenMP/target_data_use_device_addr_codegen.cpp
@@ -12,11 +12,11 @@
#define HEADER
// CHECK-DAG: [[SIZES1:@.+]] = private unnamed_addr constant [5 x i64] zeroinitializer
-// 96 = 0x60 = OMP_MAP_TARGET_PARAM | OMP_MAP_RETURN_PARAM
-// CHECK-DAG: [[MAPTYPES1:@.+]] = private unnamed_addr constant [5 x i64] [i64 96, i64 96, i64 96, i64 96, i64 96]
-// 32 = 0x20 = OMP_MAP_TARGET_PARAM
+// 64 = 0x40 = OMP_MAP_RETURN_PARAM
+// CHECK-DAG: [[MAPTYPES1:@.+]] = private unnamed_addr constant [5 x i64] [i64 64, i64 64, i64 64, i64 64, i64 64]
+// 0 = OMP_MAP_NONE
// 281474976710720 = 0x1000000000040 = OMP_MAP_MEMBER_OF | OMP_MAP_RETURN_PARAM
-// CHECK-DAG: [[MAPTYPES2:@.+]] = private unnamed_addr constant [5 x i64] [i64 32, i64 281474976710720, i64 281474976710720, i64 281474976710720, i64 281474976710720]
+// CHECK-DAG: [[MAPTYPES2:@.+]] = private unnamed_addr constant [5 x i64] [i64 0, i64 281474976710720, i64 281474976710720, i64 281474976710720, i64 281474976710720]
struct S {
int a = 0;
int *ptr = &a;
diff --git a/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp b/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp
--- a/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp
+++ b/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp
@@ -22,18 +22,18 @@
double *g;
// CK1: @g = global double*
-// CK1: [[MTYPE00:@.+]] = {{.*}}constant [2 x i64] [i64 51, i64 96]
-// CK1: [[MTYPE01:@.+]] = {{.*}}constant [1 x i64] [i64 99]
-// CK1: [[MTYPE03:@.+]] = {{.*}}constant [1 x i64] [i64 99]
-// CK1: [[MTYPE04:@.+]] = {{.*}}constant [1 x i64] [i64 99]
-// CK1: [[MTYPE05:@.+]] = {{.*}}constant [1 x i64] [i64 99]
-// CK1: [[MTYPE06:@.+]] = {{.*}}constant [1 x i64] [i64 99]
-// CK1: [[MTYPE07:@.+]] = {{.*}}constant [1 x i64] [i64 99]
-// CK1: [[MTYPE08:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 35]
-// CK1: [[MTYPE09:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 99]
-// CK1: [[MTYPE10:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 99]
-// CK1: [[MTYPE11:@.+]] = {{.*}}constant [2 x i64] [i64 35, i64 96]
-// CK1: [[MTYPE12:@.+]] = {{.*}}constant [2 x i64] [i64 35, i64 96]
+// CK1: [[MTYPE00:@.+]] = {{.*}}constant [2 x i64] [i64 19, i64 64]
+// CK1: [[MTYPE01:@.+]] = {{.*}}constant [1 x i64] [i64 67]
+// CK1: [[MTYPE03:@.+]] = {{.*}}constant [1 x i64] [i64 67]
+// CK1: [[MTYPE04:@.+]] = {{.*}}constant [1 x i64] [i64 67]
+// CK1: [[MTYPE05:@.+]] = {{.*}}constant [1 x i64] [i64 67]
+// CK1: [[MTYPE06:@.+]] = {{.*}}constant [1 x i64] [i64 67]
+// CK1: [[MTYPE07:@.+]] = {{.*}}constant [1 x i64] [i64 67]
+// CK1: [[MTYPE08:@.+]] = {{.*}}constant [2 x i64] [i64 67, i64 3]
+// CK1: [[MTYPE09:@.+]] = {{.*}}constant [2 x i64] [i64 67, i64 67]
+// CK1: [[MTYPE10:@.+]] = {{.*}}constant [2 x i64] [i64 67, i64 67]
+// CK1: [[MTYPE11:@.+]] = {{.*}}constant [2 x i64] [i64 3, i64 64]
+// CK1: [[MTYPE12:@.+]] = {{.*}}constant [2 x i64] [i64 3, i64 64]
// CK1-LABEL: @_Z3foo
template
@@ -346,10 +346,10 @@
#ifdef CK2
// CK2: [[ST:%.+]] = type { double*, double** }
-// CK2: [[MTYPE00:@.+]] = {{.*}}constant [2 x i64] [i64 32, i64 281474976710739]
-// CK2: [[MTYPE01:@.+]] = {{.*}}constant [2 x i64] [i64 32, i64 281474976710739]
-// CK2: [[MTYPE02:@.+]] = {{.*}}constant [3 x i64] [i64 35, i64 32, i64 562949953421392]
-// CK2: [[MTYPE03:@.+]] = {{.*}}constant [3 x i64] [i64 32, i64 281474976710739, i64 281474976710736]
+// CK2: [[MTYPE00:@.+]] = {{.*}}constant [2 x i64] [i64 0, i64 281474976710739]
+// CK2: [[MTYPE01:@.+]] = {{.*}}constant [2 x i64] [i64 0, i64 281474976710739]
+// CK2: [[MTYPE02:@.+]] = {{.*}}constant [3 x i64] [i64 3, i64 0, i64 562949953421392]
+// CK2: [[MTYPE03:@.+]] = {{.*}}constant [3 x i64] [i64 0, i64 281474976710739, i64 281474976710736]
template
struct ST {
diff --git a/clang/test/OpenMP/target_enter_data_codegen.cpp b/clang/test/OpenMP/target_enter_data_codegen.cpp
--- a/clang/test/OpenMP/target_enter_data_codegen.cpp
+++ b/clang/test/OpenMP/target_enter_data_codegen.cpp
@@ -30,19 +30,19 @@
double gc[100];
// CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 800]
-// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 32]
+// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] zeroinitializer
// CK1: [[SIZE02:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 4]
-// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 1]
-// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 37]
+// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 5]
// CK1: [[SIZE04:@.+]] = {{.+}}constant [2 x i64] [i64 sdiv exact (i64 sub (i64 ptrtoint (double** getelementptr (double*, double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1), i32 1) to i64), i64 ptrtoint (double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1) to i64)), i64 ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)), i64 24]
-// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673]
+// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710673]
-// CK1: [[MTYPE05:@.+]] = {{.+}}constant [1 x i64] [i64 1057]
+// CK1: [[MTYPE05:@.+]] = {{.+}}constant [1 x i64] [i64 1025]
-// CK1: [[MTYPE06:@.+]] = {{.+}}constant [1 x i64] [i64 1061]
+// CK1: [[MTYPE06:@.+]] = {{.+}}constant [1 x i64] [i64 1029]
// CK1-LABEL: _Z3fooi
void foo(int arg) {
@@ -229,11 +229,11 @@
ST gb;
double gc[100];
-// PRESENT=0x1000 | TARGET_PARAM=0x20 | TO=0x1 = 0x1021
-// CK1A: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1021]]]
+// PRESENT=0x1000 | TO=0x1 = 0x1001
+// CK1A: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1001]]]
-// PRESENT=0x1000 | CLOSE=0x400 | TARGET_PARAM=0x20 | ALWAYS=0x4 | TO=0x1 = 0x1425
-// CK1A: [[MTYPE01:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1425]]]
+// PRESENT=0x1000 | CLOSE=0x400 | ALWAYS=0x4 | TO=0x1 = 0x1405
+// CK1A: [[MTYPE01:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1405]]]
// CK1A-LABEL: _Z3fooi
void foo(int arg) {
@@ -319,7 +319,7 @@
}
};
-// CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710677]
+// CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710677]
// CK2-LABEL: _Z3bari
int bar(int arg){
@@ -469,7 +469,7 @@
}
};
-// CK5: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976711701]
+// CK5: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976711701]
// CK5-LABEL: _Z3bari
int bar(int arg){
diff --git a/clang/test/OpenMP/target_enter_data_depend_codegen.cpp b/clang/test/OpenMP/target_enter_data_depend_codegen.cpp
--- a/clang/test/OpenMP/target_enter_data_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_enter_data_depend_codegen.cpp
@@ -30,15 +30,15 @@
double gc[100];
// CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i64] [i64 800]
-// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 32]
+// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] zeroinitializer
// CK1: [[SIZE02:@.+]] = {{.+}}constant [1 x i64] [i64 4]
-// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 1]
-// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 32]
+// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] zeroinitializer
// CK1: [[SIZE04:@.+]] = {{.+}}constant [2 x i64] [i64 sdiv exact (i64 sub (i64 ptrtoint (double** getelementptr (double*, double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1), i32 1) to i64), i64 ptrtoint (double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1) to i64)), i64 ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)), i64 24]
-// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673]
+// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710673]
// CK1-LABEL: _Z3fooi
void foo(int arg) {
diff --git a/clang/test/OpenMP/target_exit_data_codegen.cpp b/clang/test/OpenMP/target_exit_data_codegen.cpp
--- a/clang/test/OpenMP/target_exit_data_codegen.cpp
+++ b/clang/test/OpenMP/target_exit_data_codegen.cpp
@@ -30,19 +30,19 @@
double gc[100];
// CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 800]
-// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 34]
+// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 2]
// CK1: [[SIZE02:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 4]
-// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 32]
+// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] zeroinitializer
-// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 38]
+// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 6]
// CK1: [[SIZE04:@.+]] = {{.+}}constant [2 x i64] [i64 sdiv exact (i64 sub (i64 ptrtoint (double** getelementptr (double*, double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1), i32 1) to i64), i64 ptrtoint (double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1) to i64)), i64 ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)), i64 24]
-// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710672]
+// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710672]
-// CK1: [[MTYPE05:@.+]] = {{.+}}constant [1 x i64] [i64 1058]
+// CK1: [[MTYPE05:@.+]] = {{.+}}constant [1 x i64] [i64 1026]
-// CK1: [[MTYPE06:@.+]] = {{.+}}constant [1 x i64] [i64 1062]
+// CK1: [[MTYPE06:@.+]] = {{.+}}constant [1 x i64] [i64 1030]
// CK1-LABEL: _Z3fooi
void foo(int arg) {
@@ -233,7 +233,7 @@
}
};
-// CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710676]
+// CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710676]
// CK2-LABEL: _Z3bari
int bar(int arg){
@@ -337,7 +337,7 @@
}
};
-// CK4: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976711700]
+// CK4: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976711700]
// CK4-LABEL: _Z3bari
int bar(int arg){
diff --git a/clang/test/OpenMP/target_exit_data_depend_codegen.cpp b/clang/test/OpenMP/target_exit_data_depend_codegen.cpp
--- a/clang/test/OpenMP/target_exit_data_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_exit_data_depend_codegen.cpp
@@ -30,15 +30,15 @@
double gc[100];
// CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i64] [i64 800]
-// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 34]
+// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 2]
// CK1: [[SIZE02:@.+]] = {{.+}}constant [1 x i64] [i64 4]
-// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 40]
+// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 8]
-// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 34]
+// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 2]
// CK1: [[SIZE04:@.+]] = {{.+}}constant [2 x i64] [i64 sdiv exact (i64 sub (i64 ptrtoint (double** getelementptr (double*, double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1), i32 1) to i64), i64 ptrtoint (double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1) to i64)), i64 ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)), i64 24]
-// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710674]
+// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710674]
// CK1-LABEL: _Z3fooi
void foo(int arg) {
diff --git a/clang/test/OpenMP/target_is_device_ptr_messages.cpp b/clang/test/OpenMP/target_is_device_ptr_messages.cpp
--- a/clang/test/OpenMP/target_is_device_ptr_messages.cpp
+++ b/clang/test/OpenMP/target_is_device_ptr_messages.cpp
@@ -1,6 +1,8 @@
-// RUN: %clang_cc1 -std=c++11 -verify -fopenmp -ferror-limit 200 %s -Wuninitialized
+// RUN: %clang_cc1 -std=c++11 -verify=expected,omp45 -fopenmp -fopenmp-version=45 -ferror-limit 200 %s -Wuninitialized
+// RUN: %clang_cc1 -std=c++11 -verify=expected,omp50 -fopenmp -ferror-limit 200 %s -Wuninitialized
-// RUN: %clang_cc1 -std=c++11 -verify -fopenmp-simd -ferror-limit 200 %s -Wuninitialized
+// RUN: %clang_cc1 -std=c++11 -verify=expected,omp45 -fopenmp-simd -fopenmp-version=45 -ferror-limit 200 %s -Wuninitialized
+// RUN: %clang_cc1 -std=c++11 -verify=expected,omp50 -fopenmp-simd -ferror-limit 200 %s -Wuninitialized
struct ST {
int *a;
};
diff --git a/clang/test/OpenMP/target_map_codegen_29.cpp b/clang/test/OpenMP/target_map_codegen_29.cpp
--- a/clang/test/OpenMP/target_map_codegen_29.cpp
+++ b/clang/test/OpenMP/target_map_codegen_29.cpp
@@ -73,9 +73,8 @@
// CK30-DAG: [[S_BEGIN_BC]] = ptrtoint i8* [[S_BEGIN:%.+]] to i64
// CK30-DAG: [[S_END_BC]] = ptrtoint i8* [[S_END:%.+]] to i64
// CK30-DAG: [[S_BEGIN]] = bitcast [[STRUCT]]* [[S]] to i8*
-// CK30-DAG: [[S_END]] = getelementptr i8, i8* [[S_LAST:%.+]], i32 1
-// CK30-DAG: [[S_LAST]] = getelementptr i8, i8* [[S_BC:%.+]], i{{64|32}} {{55|27}}
-// CK30-DAG: [[S_BC]] = bitcast [[STRUCT]]* [[S]] to i8*
+// CK30-DAG: [[S_END]] = bitcast [[STRUCT]]* [[REAL_S_END:%.+]] to i8*
+// CK30-DAG: [[REAL_S_END]] = getelementptr [[STRUCT]], [[STRUCT]]* [[S]], i32 1
// CK30-DAG: [[BASE_PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BASES]], i32 0, i32 1
// CK30-DAG: [[BC:%.+]] = bitcast i8** [[BASE_PTR]] to [[STRUCT]]**
@@ -125,7 +124,9 @@
// CK30-DAG: [[S_PTR1_BC]] = ptrtoint i8* [[S_PTR1:%.+]] to i64
// CK30-DAG: [[S_END_BC]] = ptrtoint i8* [[S_END:%.+]] to i64
// CK30-DAG: [[S_PTR1]] = bitcast i32** [[PTR2]] to i8*
-// CK30-DAG: [[S_END]] = getelementptr i8, i8* [[S_LAST]], i{{64|32}} 1
+// CK30-DAG: [[S_END]] = getelementptr i8, i8* [[S_LAST:%.+]], i{{64|32}} 1
+// CK30-DAG: [[S_LAST]] = getelementptr i8, i8* [[S_BC:%.+]], i{{64|32}} {{55|27}}
+// CK30-DAG: [[S_BC]] = bitcast [[STRUCT]]* [[S]] to i8*
// CK30-DAG: [[BASE_PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BASES]], i32 0, i32 4
// CK30-DAG: [[BC:%.+]] = bitcast i8** [[BASE_PTR]] to i32***
diff --git a/clang/test/OpenMP/target_map_codegen_31.cpp b/clang/test/OpenMP/target_map_codegen_31.cpp
--- a/clang/test/OpenMP/target_map_codegen_31.cpp
+++ b/clang/test/OpenMP/target_map_codegen_31.cpp
@@ -45,11 +45,11 @@
// CK31A-USE: [[MTYPE00:@.+]] = private {{.*}}constant [7 x i64] [i64 [[#0x1020]],
// CK31A-NOUSE: [[MTYPE00:@.+]] = private {{.*}}constant [7 x i64] [i64 [[#0x1000]],
//
-// MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003
// MEMBER_OF_1=0x1000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x1000000001003
+// MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003
// PRESENT=0x1000 | TARGET_PARAM=0x20 | FROM=0x2 | TO=0x1 = 0x1023
-// CK31A-USE-SAME: {{^}} i64 [[#0x1000000000003]], i64 [[#0x1000000001003]], i64 [[#0x1023]],
-// CK31A-NOUSE-SAME: {{^}} i64 [[#0x1000000000003]], i64 [[#0x1000000001003]], i64 [[#0x1003]],
+// CK31A-USE-SAME: {{^}} i64 [[#0x1000000001003]], i64 [[#0x1000000000003]], i64 [[#0x1023]],
+// CK31A-NOUSE-SAME: {{^}} i64 [[#0x1000000001003]], i64 [[#0x1000000000003]], i64 [[#0x1003]],
//
// PRESENT=0x1000 | TARGET_PARAM=0x20 = 0x1020
// MEMBER_OF_5=0x5000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x5000000001003
@@ -77,91 +77,94 @@
struct ST st1;
struct ST st2;
- // Make sure the struct picks up present even if another element of the struct
- // doesn't have present.
- // Region 00
- // CK31A: [[ST1_I:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST1]], i{{.+}} 0, i{{.+}} 0
- // CK31A: [[ST1_J:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST1]], i{{.+}} 0, i{{.+}} 1
- // CK31A: [[ST2_I:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST2]], i{{.+}} 0, i{{.+}} 0
- // CK31A: [[ST2_J:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST2]], i{{.+}} 0, i{{.+}} 1
- // CK31A-DAG: call i32 @__tgt_target_mapper(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(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 7, i8** [[GEPBP:%[0-9]+]], i8** [[GEPP:%[0-9]+]], i64* [[GEPS:%.+]], i64* getelementptr {{.+}}[7 x i{{.+}}]* [[MTYPE00]]{{.+}})
+// CK31A-DAG: [[GEPS]] = getelementptr inbounds [7 x i64], [7 x i64]* [[S:%.+]], i{{.+}} 0, i{{.+}} 0
+// CK31A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+// CK31A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+// st1
+// CK31A-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+// CK31A-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+// CK31A-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+// CK31A-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
+// CK31A-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
+// CK31A-DAG: store [[ST]]* [[ST1]], [[ST]]** [[CBP0]]
+// CK31A-DAG: store i32* [[ST1_I]], i32** [[CP0]]
+// CK31A-DAG: store i64 %{{.+}}, i64* [[S0]]
+
+// st1.j
+// CK31A-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+// CK31A-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+// CK31A-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
+// CK31A-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]**
+// CK31A-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32**
+// CK31A-DAG: store [[ST]]* [[ST1]], [[ST]]** [[CBP1]]
+// CK31A-DAG: store i32* [[ST1_J]], i32** [[CP1]]
+// CK31A-DAG: store i64 4, i64* [[S1]]
+
+// st1.i
+// CK31A-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
+// CK31A-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
+// CK31A-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2
+// CK31A-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to [[ST]]**
+// CK31A-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to i32**
+// CK31A-DAG: store [[ST]]* [[ST1]], [[ST]]** [[CBP2]]
+// CK31A-DAG: store i32* [[ST1_I]], i32** [[CP2]]
+// CK31A-DAG: store i64 4, i64* [[S2]]
+
+// a
+// CK31A-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3
+// CK31A-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3
+// CK31A-DAG: [[S3:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 3
+// CK31A-DAG: [[CBP3:%.+]] = bitcast i8** [[BP3]] to i32**
+// CK31A-DAG: [[CP3:%.+]] = bitcast i8** [[P3]] to i32**
+// CK31A-DAG: store i32* [[A]], i32** [[CBP3]]
+// CK31A-DAG: store i32* [[A]], i32** [[CP3]]
+// CK31A-DAG: store i64 4, i64* [[S3]]
+
+// st2
+// CK31A-DAG: [[BP4:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 4
+// CK31A-DAG: [[P4:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 4
+// CK31A-DAG: [[S4:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 4
+// CK31A-DAG: [[CBP4:%.+]] = bitcast i8** [[BP4]] to [[ST]]**
+// CK31A-DAG: [[CP4:%.+]] = bitcast i8** [[P4]] to i32**
+// CK31A-DAG: store [[ST]]* [[ST2]], [[ST]]** [[CBP4]]
+// CK31A-DAG: store i32* [[ST2_I]], i32** [[CP4]]
+// CK31A-DAG: store i64 %{{.+}}, i64* [[S4]]
+
+// st2.i
+// CK31A-DAG: [[BP5:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 5
+// CK31A-DAG: [[P5:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 5
+// CK31A-DAG: [[S5:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 5
+// CK31A-DAG: [[CBP5:%.+]] = bitcast i8** [[BP5]] to [[ST]]**
+// CK31A-DAG: [[CP5:%.+]] = bitcast i8** [[P5]] to i32**
+// CK31A-DAG: store [[ST]]* [[ST2]], [[ST]]** [[CBP5]]
+// CK31A-DAG: store i32* [[ST2_I]], i32** [[CP5]]
+// CK31A-DAG: store i64 4, i64* [[S5]]
+
+// st2.j
+// CK31A-DAG: [[BP6:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 6
+// CK31A-DAG: [[P6:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 6
+// CK31A-DAG: [[S6:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 6
+// CK31A-DAG: [[CBP6:%.+]] = bitcast i8** [[BP6]] to [[ST]]**
+// CK31A-DAG: [[CP6:%.+]] = bitcast i8** [[P6]] to i32**
+// CK31A-DAG: store [[ST]]* [[ST2]], [[ST]]** [[CBP6]]
+// CK31A-DAG: store i32* [[ST2_J]], i32** [[CP6]]
+// CK31A-DAG: store i64 4, i64* [[S6]]
+
+// CK31A-USE: call void [[CALL00:@.+]]([[ST]]* [[ST1]], i32* [[A]], [[ST]]* [[ST2]])
+// CK31A-NOUSE: call void [[CALL00:@.+]]()
+#pragma omp target map(tofrom \
+ : st1.i) map(present, tofrom \
+ : a, st1.j, st2.i) map(tofrom \
+ : st2.j)
{
#ifdef USE
st1.i++;
diff --git a/clang/test/OpenMP/target_map_codegen_32.cpp b/clang/test/OpenMP/target_map_codegen_32.cpp
--- a/clang/test/OpenMP/target_map_codegen_32.cpp
+++ b/clang/test/OpenMP/target_map_codegen_32.cpp
@@ -36,63 +36,65 @@
// CK31B: [[ST:%.+]] = type { i32, i32 }
// PRESENT=0x1000 | TARGET_PARAM=0x20 = 0x1020
-// MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003
// MEMBER_OF_1=0x1000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x1000000001003
+// MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003
// CK31B-LABEL: @.__omp_offloading_{{.*}}test_present_members{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
// CK31B-USE: [[MTYPE00:@.+]] = private {{.*}}constant [3 x i64] [i64 [[#0x1020]],
// CK31B-NOUSE: [[MTYPE00:@.+]] = private {{.*}}constant [3 x i64] [i64 [[#0x1000]],
-// CK31B-USE-SAME: {{^}} i64 [[#0x1000000000003]], i64 [[#0x1000000001003]]]
-// CK31B-NOUSE-SAME: {{^}} i64 [[#0x1000000000003]], i64 [[#0x1000000001003]]]
+// CK31B-USE-SAME: {{^}} i64 [[#0x1000000001003]], i64 [[#0x1000000000003]]]
+// CK31B-NOUSE-SAME: {{^}} i64 [[#0x1000000001003]], i64 [[#0x1000000000003]]]
struct ST {
int i;
int j;
// CK31B-LABEL: define {{.*}}test_present_members{{.*}}(
void test_present_members() {
- // Make sure the struct picks up present even if another element of the
- // struct doesn't have present.
- // Region 00
- // CK31B: [[I:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[THIS:%.+]], i{{.+}} 0, i{{.+}} 0
- // CK31B: [[J:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[THIS]], i{{.+}} 0, i{{.+}} 1
- // CK31B-DAG: call i32 @__tgt_target_mapper(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(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_member_expr_array_section_codegen.cpp b/clang/test/OpenMP/target_map_member_expr_array_section_codegen.cpp
--- a/clang/test/OpenMP/target_map_member_expr_array_section_codegen.cpp
+++ b/clang/test/OpenMP/target_map_member_expr_array_section_codegen.cpp
@@ -5,11 +5,11 @@
#ifndef HEADER
#define HEADER
-// 32 = 0x20 = OMP_MAP_TARGET_PARAM
+// 0 = OMP_MAP_NONE
// 281474976710656 = 0x1000000000000 = OMP_MAP_MEMBER_OF of 1-st element
-// CHECK: [[MAP_ENTER:@.+]] = private unnamed_addr constant [2 x i64] [i64 32, i64 281474976710656]
+// CHECK: [[MAP_ENTER:@.+]] = private unnamed_addr constant [2 x i64] [i64 0, i64 281474976710656]
// 281474976710664 = 0x1000000000008 = OMP_MAP_MEMBER_OF of 1-st element | OMP_MAP_DELETE
-// CHECK: [[MAP_EXIT:@.+]] = private unnamed_addr constant [2 x i64] [i64 32, i64 281474976710664]
+// CHECK: [[MAP_EXIT:@.+]] = private unnamed_addr constant [2 x i64] [i64 0, i64 281474976710664]
template
struct S {
constexpr static int size = 6;
diff --git a/clang/test/OpenMP/target_map_messages.cpp b/clang/test/OpenMP/target_map_messages.cpp
--- a/clang/test/OpenMP/target_map_messages.cpp
+++ b/clang/test/OpenMP/target_map_messages.cpp
@@ -79,7 +79,7 @@
#pragma omp target map(to:b,e)
{}
- #pragma omp target map(to:b,e) map(to:b) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+ #pragma omp target map(to:b,e) map(to:b) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
{}
#pragma omp target map(to:b[:2],e)
{}
@@ -338,19 +338,19 @@
{}
#pragma omp target map(r.ArrS[0].Error) // expected-error {{no member named 'Error' in 'SB'}}
{}
-#pragma omp target map(r.ArrS[0].A, r.ArrS[1].A) // expected-error {{multiple array elements associated with the same variable are not allowed in map clauses of the same construct}} expected-note {{used here}}
+#pragma omp target map(r.ArrS[0].A, r.ArrS[1].A) // lt50-error {{multiple array elements associated with the same variable are not allowed in map clauses of the same construct}} lt50-note {{used here}}
{}
#pragma omp target map(r.ArrS[0].A, t.ArrS[1].A)
{}
-#pragma omp target map(r.PtrS[0], r.PtrS->B) // expected-error {{same pointer dereferenced in multiple different ways in map clause expressions}} expected-note {{used here}}
+#pragma omp target map(r.PtrS[0], r.PtrS->B) // lt50-error {{same pointer dereferenced in multiple different ways in map clause expressions}} lt50-note {{used here}}
{}
-#pragma omp target map(r.PtrS, r.PtrS->B) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
+#pragma omp target map(r.PtrS, r.PtrS->B) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}}
{}
#pragma omp target map(r.PtrS->A, r.PtrS->B)
{}
-#pragma omp target map(r.RPtrS[0], r.RPtrS->B) // expected-error {{same pointer dereferenced in multiple different ways in map clause expressions}} expected-note {{used here}}
+#pragma omp target map(r.RPtrS[0], r.RPtrS->B) // lt50-error {{same pointer dereferenced in multiple different ways in map clause expressions}} lt50-note {{used here}}
{}
-#pragma omp target map(r.RPtrS, r.RPtrS->B) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
+#pragma omp target map(r.RPtrS, r.RPtrS->B) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}}
{}
#pragma omp target map(r.RPtrS->A, r.RPtrS->B)
{}
@@ -360,13 +360,13 @@
{}
#pragma omp target map(r.C, r.D)
{}
-#pragma omp target map(r.C, r.C) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+#pragma omp target map(r.C, r.C) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
{}
-#pragma omp target map(r.C) map(r.C) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+#pragma omp target map(r.C) map(r.C) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
{}
#pragma omp target map(r.C, r.S) // this would be an error only caught at runtime - Sema would have to make sure there is not way for the missing data between fields to be mapped somewhere else.
{}
-#pragma omp target map(r, r.S) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+#pragma omp target map(r, r.S) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
{}
#pragma omp target map(r.C, t.C)
{}
@@ -401,16 +401,16 @@
}
#pragma omp target data map(to \
- : r.C) //expected-note {{used here}}
+ : r.C) // lt50-note {{used here}}
{
-#pragma omp target map(r.D) // expected-error {{original storage of expression in data environment is shared but data environment do not fully contain mapped expression storage}}
+#pragma omp target map(r.D) // lt50-error {{original storage of expression in data environment is shared but data environment do not fully contain mapped expression storage}}
{}
}
#pragma omp target data map(to \
- : t.Ptr) //expected-note {{used here}}
+ : t.Ptr) // lt50-note {{used here}}
{
-#pragma omp target map(t.Ptr[:23]) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target map(t.Ptr[:23]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
{}
}
@@ -581,19 +581,19 @@
#pragma omp target data map(S2::S2sc)
#pragma omp target data map(e, g) // warn-warning 2 {{Type 'S4' is not trivially copyable and not guaranteed to be mapped correctly}} warn-warning 2 {{Type 'S5' is not trivially copyable and not guaranteed to be mapped correctly}}
#pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
-#pragma omp target data map(k) map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
-#pragma omp target map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}}
+#pragma omp target data map(k) map(k) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
+#pragma omp target map(k), map(k[:5]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} lt50-note 2 {{used here}}
foo();
#pragma omp target data map(da)
#pragma omp target map(da[:4])
foo();
-#pragma omp target data map(k, j, l) // expected-note 2 {{used here}}
-#pragma omp target data map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l) // lt50-note 2 {{used here}}
+#pragma omp target data map(k[:4]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
#pragma omp target data map(j)
-#pragma omp target map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+#pragma omp target map(l) map(l[:5]) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}}
-#pragma omp target data map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note 2 {{used here}}
+#pragma omp target data map(k) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
#pragma omp target data map(j)
#pragma omp target map(l)
foo();
@@ -716,19 +716,19 @@
#pragma omp target data map(S2::S2sc)
#pragma omp target data map(e, g) // warn-warning {{Type 'S4' is not trivially copyable and not guaranteed to be mapped correctly}} warn-warning {{Type 'S5' is not trivially copyable and not guaranteed to be mapped correctly}}
#pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
-#pragma omp target data map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
-#pragma omp target map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
+#pragma omp target data map(k), map(k) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
+#pragma omp target map(k), map(k[:5]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}}
foo();
#pragma omp target data map(da)
#pragma omp target map(da[:4])
foo();
-#pragma omp target data map(k, j, l) // expected-note {{used here}}
-#pragma omp target data map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l) // lt50-note {{used here}}
+#pragma omp target data map(k[:4]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
#pragma omp target data map(j)
-#pragma omp target map(l) map(l[:5]) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+#pragma omp target map(l) map(l[:5]) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note {{used here}}
-#pragma omp target data map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note {{used here}}
+#pragma omp target data map(k) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
#pragma omp target data map(j)
#pragma omp target map(l)
foo();
@@ -766,47 +766,49 @@
{}
#pragma omp target map(m) // warn-warning {{Type 'S6' is not trivially copyable and not guaranteed to be mapped correctly}}
{}
-// expected-note@+1 {{used here}}
+#pragma omp target
+ { s.a++; }
+// lt50-note@+1 {{used here}}
#pragma omp target map(s.s.s)
-// expected-error@+1 {{variable already marked as mapped in current construct}}
+// lt50-error@+1 {{variable already marked as mapped in current construct}}
{ s.a++; }
-// expected-note@+1 {{used here}}
+// lt50-note@+1 {{used here}}
#pragma omp target map(s.s.s.a)
-// expected-error@+1 {{variable already marked as mapped in current construct}}
+// lt50-error@+1 {{variable already marked as mapped in current construct}}
{ s.a++; }
-// expected-note@+1 {{used here}}
+// lt50-note@+1 {{used here}}
#pragma omp target map(s.b[:5])
-// expected-error@+1 {{variable already marked as mapped in current construct}}
+// lt50-error@+1 {{variable already marked as mapped in current construct}}
{ s.a++; }
#pragma omp target map(s.p[:5])
{ s.a++; }
-// expected-note@+1 {{used here}}
+// lt50-note@+1 {{used here}}
#pragma omp target map(s.s.sa[3].a)
-// expected-error@+1 {{variable already marked as mapped in current construct}}
+// lt50-error@+1 {{variable already marked as mapped in current construct}}
{ s.a++; }
-// expected-note@+1 {{used here}}
+// lt50-note@+1 {{used here}}
#pragma omp target map(s.s.sp[3]->a)
-// expected-error@+1 {{variable already marked as mapped in current construct}}
+// lt50-error@+1 {{variable already marked as mapped in current construct}}
{ s.a++; }
-// expected-note@+1 {{used here}}
+// lt50-note@+1 {{used here}}
#pragma omp target map(s.p->a)
-// expected-error@+1 {{variable already marked as mapped in current construct}}
+// lt50-error@+1 {{variable already marked as mapped in current construct}}
{ s.a++; }
-// expected-note@+1 {{used here}}
+// lt50-note@+1 {{used here}}
#pragma omp target map(s.s.p->a)
-// expected-error@+1 {{variable already marked as mapped in current construct}}
+// lt50-error@+1 {{variable already marked as mapped in current construct}}
{ s.a++; }
-// expected-note@+1 {{used here}}
+// lt50-note@+1 {{used here}}
#pragma omp target map(s.s.s.b[:2])
-// expected-error@+1 {{variable already marked as mapped in current construct}}
+// lt50-error@+1 {{variable already marked as mapped in current construct}}
{ s.a++; }
-// expected-note@+1 {{used here}}
+// lt50-note@+1 {{used here}}
#pragma omp target map(s.s.p->b[:2])
-// expected-error@+1 {{variable already marked as mapped in current construct}}
+// lt50-error@+1 {{variable already marked as mapped in current construct}}
{ s.a++; }
-// expected-note@+1 {{used here}}
+// lt50-note@+1 {{used here}}
#pragma omp target map(s.p->p->p->a)
-// expected-error@+1 {{variable already marked as mapped in current construct}}
+// lt50-error@+1 {{variable already marked as mapped in current construct}}
{ s.a++; }
#pragma omp target map(s.s.s.b[:2])
{ s.s.s.b[0]++; }
diff --git a/clang/test/OpenMP/target_parallel_for_map_messages.cpp b/clang/test/OpenMP/target_parallel_for_map_messages.cpp
--- a/clang/test/OpenMP/target_parallel_for_map_messages.cpp
+++ b/clang/test/OpenMP/target_parallel_for_map_messages.cpp
@@ -152,25 +152,26 @@
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+#pragma omp target parallel for map(k), map(k) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}}
+#pragma omp target parallel for map(k), map(k[:5]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} lt50-note 2 {{used here}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for map(da)
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for map(da[:4])
for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k, j, l) // expected-note 2 {{used here}}
-#pragma omp target parallel for map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l) // lt50-note 2 {{used here}}
+#pragma omp target parallel for map(k[:4]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for map(j)
for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
- for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}}
-{
-#pragma omp target parallel for map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target parallel for map(l) map(l[:5]) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
for (i = 0; i < argc; ++i) foo();
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note 2 {{used here}}
+ {
+#pragma omp target parallel for map(k) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+ for (i = 0; i < argc; ++i)
+ foo();
#pragma omp target parallel for map(j)
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for map(l)
@@ -269,25 +270,26 @@
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+#pragma omp target parallel for map(k), map(k) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
+#pragma omp target parallel for map(k), map(k[:5]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for map(da)
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for map(da[:4])
for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k, j, l) // expected-note {{used here}}
-#pragma omp target parallel for map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l) // lt50-note {{used here}}
+#pragma omp target parallel for map(k[:4]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for map(j)
for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for map(l) map(l[:5]) // expected-error 1 {{variable already marked as mapped in current construct}} expected-note 1 {{used here}}
- for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note {{used here}}
-{
-#pragma omp target parallel for map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target parallel for map(l) map(l[:5]) // lt50-error 1 {{variable already marked as mapped in current construct}} lt50-note 1 {{used here}}
for (i = 0; i < argc; ++i) foo();
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note {{used here}}
+ {
+#pragma omp target parallel for map(k) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
+ for (i = 0; i < argc; ++i)
+ foo();
#pragma omp target parallel for map(j)
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for map(l)
diff --git a/clang/test/OpenMP/target_parallel_for_simd_map_messages.cpp b/clang/test/OpenMP/target_parallel_for_simd_map_messages.cpp
--- a/clang/test/OpenMP/target_parallel_for_simd_map_messages.cpp
+++ b/clang/test/OpenMP/target_parallel_for_simd_map_messages.cpp
@@ -152,25 +152,26 @@
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for simd map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for simd map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+#pragma omp target parallel for simd map(k), map(k) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for simd map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}}
+#pragma omp target parallel for simd map(k), map(k[:5]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} lt50-note 2 {{used here}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for simd map(da)
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for simd map(da[:4])
for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k, j, l) // expected-note 2 {{used here}}
-#pragma omp target parallel for simd map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l) // lt50-note 2 {{used here}}
+#pragma omp target parallel for simd map(k[:4]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for simd map(j)
for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for simd map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
- for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}}
-{
-#pragma omp target parallel for simd map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target parallel for simd map(l) map(l[:5]) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
for (i = 0; i < argc; ++i) foo();
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note 2 {{used here}}
+ {
+#pragma omp target parallel for simd map(k) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+ for (i = 0; i < argc; ++i)
+ foo();
#pragma omp target parallel for simd map(j)
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for simd map(l)
@@ -269,25 +270,26 @@
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for simd map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for simd map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+#pragma omp target parallel for simd map(k), map(k) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for simd map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
+#pragma omp target parallel for simd map(k), map(k[:5]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for simd map(da)
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for simd map(da[:4])
for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k, j, l) // expected-note {{used here}}
-#pragma omp target parallel for simd map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l) // lt50-note {{used here}}
+#pragma omp target parallel for simd map(k[:4]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for simd map(j)
for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for simd map(l) map(l[:5]) // expected-error 1 {{variable already marked as mapped in current construct}} expected-note 1 {{used here}}
- for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note {{used here}}
-{
-#pragma omp target parallel for simd map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target parallel for simd map(l) map(l[:5]) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
for (i = 0; i < argc; ++i) foo();
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note {{used here}}
+ {
+#pragma omp target parallel for simd map(k) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
+ for (i = 0; i < argc; ++i)
+ foo();
#pragma omp target parallel for simd map(j)
for (i = 0; i < argc; ++i) foo();
#pragma omp target parallel for simd map(l)
diff --git a/clang/test/OpenMP/target_parallel_map_messages.cpp b/clang/test/OpenMP/target_parallel_map_messages.cpp
--- a/clang/test/OpenMP/target_parallel_map_messages.cpp
+++ b/clang/test/OpenMP/target_parallel_map_messages.cpp
@@ -152,25 +152,25 @@
foo();
#pragma omp target parallel map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
foo();
-#pragma omp target parallel map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+#pragma omp target parallel map(k), map(k) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
foo();
-#pragma omp target parallel map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}}
+#pragma omp target parallel map(k), map(k[:5]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} lt50-note 2 {{used here}}
foo();
#pragma omp target parallel map(da)
foo();
#pragma omp target parallel map(da[:4])
foo();
-#pragma omp target data map(k, j, l) // expected-note 2 {{used here}}
-#pragma omp target parallel map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l) // lt50-note 2 {{used here}}
+#pragma omp target parallel map(k[:4]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
foo();
#pragma omp target parallel map(j)
foo();
-#pragma omp target parallel map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
- foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}}
-{
-#pragma omp target parallel map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target parallel map(l) map(l[:5]) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
foo();
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note 2 {{used here}}
+ {
+#pragma omp target parallel map(k) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+ foo();
#pragma omp target parallel map(j)
foo();
#pragma omp target parallel map(l)
@@ -267,25 +267,25 @@
foo();
#pragma omp target parallel map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
foo();
-#pragma omp target parallel map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+#pragma omp target parallel map(k), map(k) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
foo();
-#pragma omp target parallel map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
+#pragma omp target parallel map(k), map(k[:5]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}}
foo();
#pragma omp target parallel map(da)
foo();
#pragma omp target parallel map(da[:4])
foo();
-#pragma omp target data map(k, j, l) // expected-note {{used here}}
-#pragma omp target parallel map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l) // lt50-note {{used here}}
+#pragma omp target parallel map(k[:4]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
foo();
#pragma omp target parallel map(j)
foo();
-#pragma omp target parallel map(l) map(l[:5]) // expected-error 1 {{variable already marked as mapped in current construct}} expected-note 1 {{used here}}
- foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note 1 {{used here}}
-{
-#pragma omp target parallel map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target parallel map(l) map(l[:5]) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
foo();
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note {{used here}}
+ {
+#pragma omp target parallel map(k) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
+ foo();
#pragma omp target parallel map(j)
foo();
#pragma omp target parallel map(l)
diff --git a/clang/test/OpenMP/target_simd_map_messages.cpp b/clang/test/OpenMP/target_simd_map_messages.cpp
--- a/clang/test/OpenMP/target_simd_map_messages.cpp
+++ b/clang/test/OpenMP/target_simd_map_messages.cpp
@@ -146,25 +146,26 @@
for (i = 0; i < argc; ++i) foo();
#pragma omp target simd map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target simd map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+#pragma omp target simd map(k), map(k) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target simd map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}}
+#pragma omp target simd map(k), map(k[:5]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} lt50-note 2 {{used here}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target simd map(da)
for (i = 0; i < argc; ++i) foo();
#pragma omp target simd map(da[:4])
for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k, j, l) // expected-note 2 {{used here}}
-#pragma omp target simd map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l) // lt50-note 2 {{used here}}
+#pragma omp target simd map(k[:4]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target simd map(j)
for (i = 0; i < argc; ++i) foo();
-#pragma omp target simd map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
- for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}}
-{
-#pragma omp target simd map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target simd map(l) map(l[:5]) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
for (i = 0; i < argc; ++i) foo();
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note 2 {{used here}}
+ {
+#pragma omp target simd map(k) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+ for (i = 0; i < argc; ++i)
+ foo();
#pragma omp target simd map(j)
for (i = 0; i < argc; ++i) foo();
#pragma omp target simd map(l) // OK
@@ -256,31 +257,37 @@
#pragma omp target simd map(e, g)
for (i = 0; i < argc; ++i) foo();
#pragma omp target simd map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
- for (i = 0; i < argc; ++i) foo();
-#pragma omp target simd map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
- for (i = 0; i < argc; ++i) foo();
-#pragma omp target simd map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
+ for (i = 0; i < argc; ++i)
+ foo();
+#pragma omp target simd map(k), map(k) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
+ for (i = 0; i < argc; ++i)
+ foo();
+#pragma omp target simd map(k), map(k[:5]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target simd map(da)
for (i = 0; i < argc; ++i) foo();
#pragma omp target simd map(da[:4])
- for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k, j, l) // expected-note {{used here}}
-#pragma omp target simd map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+ for (i = 0; i < argc; ++i)
+ foo();
+#pragma omp target data map(k, j, l) // lt50-note {{used here}}
+#pragma omp target simd map(k[:4]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target simd map(j)
- for (i = 0; i < argc; ++i) foo();
-#pragma omp target simd map(l) map(l[:5]) // expected-error 1 {{variable already marked as mapped in current construct}} expected-note 1 {{used here}}
- for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note 1 {{used here}}
-{
-#pragma omp target simd map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
- for (i = 0; i < argc; ++i) foo();
+ for (i = 0; i < argc; ++i)
+ foo();
+#pragma omp target simd map(l) map(l[:5]) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
+ for (i = 0; i < argc; ++i)
+ foo();
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note {{used here}}
+ {
+#pragma omp target simd map(k) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
+ for (i = 0; i < argc; ++i)
+ foo();
#pragma omp target simd map(j)
for (i = 0; i < argc; ++i) foo();
#pragma omp target simd map(l) //
for (i = 0; i < argc; ++i) foo();
-}
+ }
#pragma omp target simd map(always, tofrom: x)
for (i = 0; i < argc; ++i) foo();
diff --git a/clang/test/OpenMP/target_teams_distribute_map_messages.cpp b/clang/test/OpenMP/target_teams_distribute_map_messages.cpp
--- a/clang/test/OpenMP/target_teams_distribute_map_messages.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_map_messages.cpp
@@ -152,25 +152,26 @@
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+#pragma omp target teams distribute map(k), map(k) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}}
+#pragma omp target teams distribute map(k), map(k[:5]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} lt50-note 2 {{used here}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute map(da)
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute map(da[:4])
for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k, j, l) // expected-note 2 {{used here}}
-#pragma omp target teams distribute map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l) // lt50-note 2 {{used here}}
+#pragma omp target teams distribute map(k[:4]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute map(j)
for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
- for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}}
-{
-#pragma omp target teams distribute map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target teams distribute map(l) map(l[:5]) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
for (i = 0; i < argc; ++i) foo();
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note 2 {{used here}}
+ {
+#pragma omp target teams distribute map(k) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+ for (i = 0; i < argc; ++i)
+ foo();
#pragma omp target teams distribute map(j)
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute map(l)
@@ -269,25 +270,26 @@
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+#pragma omp target teams distribute map(k), map(k) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
+#pragma omp target teams distribute map(k), map(k[:5]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute map(da)
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute map(da[:4])
for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k, j, l) // expected-note {{used here}}
-#pragma omp target teams distribute map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l) // lt50-note {{used here}}
+#pragma omp target teams distribute map(k[:4]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute map(j)
for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute map(l) map(l[:5]) // expected-error 1 {{variable already marked as mapped in current construct}} expected-note 1 {{used here}}
- for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note {{used here}}
-{
-#pragma omp target teams distribute map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target teams distribute map(l) map(l[:5]) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
for (i = 0; i < argc; ++i) foo();
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note {{used here}}
+ {
+#pragma omp target teams distribute map(k) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
+ for (i = 0; i < argc; ++i)
+ foo();
#pragma omp target teams distribute map(j)
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute map(l)
diff --git a/clang/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp b/clang/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp
--- a/clang/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp
@@ -150,25 +150,26 @@
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute parallel for map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+#pragma omp target teams distribute parallel for map(k), map(k) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}}
+#pragma omp target teams distribute parallel for map(k), map(k[:5]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} lt50-note 2 {{used here}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute parallel for map(da)
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute parallel for map(da[:4])
for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k, j, l) // expected-note 2 {{used here}}
-#pragma omp target teams distribute parallel for map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l) // lt50-note 2 {{used here}}
+#pragma omp target teams distribute parallel for map(k[:4]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute parallel for map(j)
for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
- for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}}
-{
-#pragma omp target teams distribute parallel for map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target teams distribute parallel for map(l) map(l[:5]) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
for (i = 0; i < argc; ++i) foo();
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note 2 {{used here}}
+ {
+#pragma omp target teams distribute parallel for map(k) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+ for (i = 0; i < argc; ++i)
+ foo();
#pragma omp target teams distribute parallel for map(j)
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute parallel for map(l)
@@ -265,31 +266,37 @@
#pragma omp target teams distribute parallel for map(e, g)
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute parallel for map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
- for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
- for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
+ for (i = 0; i < argc; ++i)
+ foo();
+#pragma omp target teams distribute parallel for map(k), map(k) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
+ for (i = 0; i < argc; ++i)
+ foo();
+#pragma omp target teams distribute parallel for map(k), map(k[:5]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute parallel for map(da)
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute parallel for map(da[:4])
- for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k, j, l) // expected-note {{used here}}
-#pragma omp target teams distribute parallel for map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+ for (i = 0; i < argc; ++i)
+ foo();
+#pragma omp target data map(k, j, l) // lt50-note {{used here}}
+#pragma omp target teams distribute parallel for map(k[:4]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute parallel for map(j)
- for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for map(l) map(l[:5]) // expected-error 1 {{variable already marked as mapped in current construct}} expected-note 1 {{used here}}
- for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note {{used here}}
-{
-#pragma omp target teams distribute parallel for map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
- for (i = 0; i < argc; ++i) foo();
+ for (i = 0; i < argc; ++i)
+ foo();
+#pragma omp target teams distribute parallel for map(l) map(l[:5]) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
+ for (i = 0; i < argc; ++i)
+ foo();
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note {{used here}}
+ {
+#pragma omp target teams distribute parallel for map(k) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
+ for (i = 0; i < argc; ++i)
+ foo();
#pragma omp target teams distribute parallel for map(j)
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute parallel for map(l)
for (i = 0; i < argc; ++i) foo();
-}
+ }
#pragma omp target teams distribute parallel for map(always, tofrom: x)
for (i = 0; i < argc; ++i) foo();
diff --git a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp
--- a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp
@@ -152,25 +152,26 @@
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute parallel for simd map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for simd map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+#pragma omp target teams distribute parallel for simd map(k), map(k) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for simd map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}}
+#pragma omp target teams distribute parallel for simd map(k), map(k[:5]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} lt50-note 2 {{used here}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute parallel for simd map(da)
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute parallel for simd map(da[:4])
for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k, j, l) // expected-note 2 {{used here}}
-#pragma omp target teams distribute parallel for simd map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l) // lt50-note 2 {{used here}}
+#pragma omp target teams distribute parallel for simd map(k[:4]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute parallel for simd map(j)
for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for simd map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
- for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}}
-{
-#pragma omp target teams distribute parallel for simd map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target teams distribute parallel for simd map(l) map(l[:5]) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
for (i = 0; i < argc; ++i) foo();
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note 2 {{used here}}
+ {
+#pragma omp target teams distribute parallel for simd map(k) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+ for (i = 0; i < argc; ++i)
+ foo();
#pragma omp target teams distribute parallel for simd map(j)
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute parallel for simd map(l)
@@ -268,25 +269,26 @@
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute parallel for simd map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for simd map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+#pragma omp target teams distribute parallel for simd map(k), map(k) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for simd map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
+#pragma omp target teams distribute parallel for simd map(k), map(k[:5]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute parallel for simd map(da)
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute parallel for simd map(da[:4])
for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k, j, l) // expected-note {{used here}}
-#pragma omp target teams distribute parallel for simd map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l) // lt50-note {{used here}}
+#pragma omp target teams distribute parallel for simd map(k[:4]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute parallel for simd map(j)
for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for simd map(l) map(l[:5]) // expected-error 1 {{variable already marked as mapped in current construct}} expected-note 1 {{used here}}
- for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note {{used here}}
-{
-#pragma omp target teams distribute parallel for simd map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target teams distribute parallel for simd map(l) map(l[:5]) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
for (i = 0; i < argc; ++i) foo();
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note {{used here}}
+ {
+#pragma omp target teams distribute parallel for simd map(k) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
+ for (i = 0; i < argc; ++i)
+ foo();
#pragma omp target teams distribute parallel for simd map(j)
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute parallel for simd map(l)
diff --git a/clang/test/OpenMP/target_teams_distribute_simd_map_messages.cpp b/clang/test/OpenMP/target_teams_distribute_simd_map_messages.cpp
--- a/clang/test/OpenMP/target_teams_distribute_simd_map_messages.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_simd_map_messages.cpp
@@ -152,25 +152,26 @@
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute simd map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute simd map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+#pragma omp target teams distribute simd map(k), map(k) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute simd map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}}
+#pragma omp target teams distribute simd map(k), map(k[:5]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} lt50-note 2 {{used here}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute simd map(da)
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute simd map(da[:4])
for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k, j, l) // expected-note 2 {{used here}}
-#pragma omp target teams distribute simd map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l) // lt50-note 2 {{used here}}
+#pragma omp target teams distribute simd map(k[:4]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute simd map(j)
for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute simd map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
- for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}}
-{
-#pragma omp target teams distribute simd map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target teams distribute simd map(l) map(l[:5]) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
for (i = 0; i < argc; ++i) foo();
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note 2 {{used here}}
+ {
+#pragma omp target teams distribute simd map(k) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+ for (i = 0; i < argc; ++i)
+ foo();
#pragma omp target teams distribute simd map(j)
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute simd map(l)
@@ -269,25 +270,26 @@
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute simd map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute simd map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+#pragma omp target teams distribute simd map(k), map(k) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute simd map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
+#pragma omp target teams distribute simd map(k), map(k[:5]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute simd map(da)
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute simd map(da[:4])
for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k, j, l) // expected-note {{used here}}
-#pragma omp target teams distribute simd map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l) // lt50-note {{used here}}
+#pragma omp target teams distribute simd map(k[:4]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute simd map(j)
for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute simd map(l) map(l[:5]) // expected-error 1 {{variable already marked as mapped in current construct}} expected-note 1 {{used here}}
- for (i = 0; i < argc; ++i) foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note {{used here}}
-{
-#pragma omp target teams distribute simd map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target teams distribute simd map(l) map(l[:5]) // lt50-error 1 {{variable already marked as mapped in current construct}} lt50-note 1 {{used here}}
for (i = 0; i < argc; ++i) foo();
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note {{used here}}
+ {
+#pragma omp target teams distribute simd map(k) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
+ for (i = 0; i < argc; ++i)
+ foo();
#pragma omp target teams distribute simd map(j)
for (i = 0; i < argc; ++i) foo();
#pragma omp target teams distribute simd map(l)
diff --git a/clang/test/OpenMP/target_teams_map_messages.cpp b/clang/test/OpenMP/target_teams_map_messages.cpp
--- a/clang/test/OpenMP/target_teams_map_messages.cpp
+++ b/clang/test/OpenMP/target_teams_map_messages.cpp
@@ -62,7 +62,7 @@
#pragma omp target teams map(to:b,e)
{}
- #pragma omp target teams map(to:b,e) map(to:b) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+ #pragma omp target teams map(to:b,e) map(to:b) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
{}
#pragma omp target teams map(to:b[:2],e)
{}
@@ -238,19 +238,19 @@
{}
#pragma omp target teams map(r.ArrS[0].Error) // expected-error {{no member named 'Error' in 'SB'}}
{}
- #pragma omp target teams map(r.ArrS[0].A, r.ArrS[1].A) // expected-error {{multiple array elements associated with the same variable are not allowed in map clauses of the same construct}} expected-note {{used here}}
+ #pragma omp target teams map(r.ArrS[0].A, r.ArrS[1].A) // lt50-error {{multiple array elements associated with the same variable are not allowed in map clauses of the same construct}} lt50-note {{used here}}
{}
#pragma omp target teams map(r.ArrS[0].A, t.ArrS[1].A)
{}
- #pragma omp target teams map(r.PtrS[0], r.PtrS->B) // expected-error {{same pointer dereferenced in multiple different ways in map clause expressions}} expected-note {{used here}}
+ #pragma omp target teams map(r.PtrS[0], r.PtrS->B) // lt50-error {{same pointer dereferenced in multiple different ways in map clause expressions}} lt50-note {{used here}}
{}
- #pragma omp target teams map(r.PtrS, r.PtrS->B) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
+ #pragma omp target teams map(r.PtrS, r.PtrS->B) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}}
{}
#pragma omp target teams map(r.PtrS->A, r.PtrS->B)
{}
- #pragma omp target teams map(r.RPtrS[0], r.RPtrS->B) // expected-error {{same pointer dereferenced in multiple different ways in map clause expressions}} expected-note {{used here}}
+ #pragma omp target teams map(r.RPtrS[0], r.RPtrS->B) // lt50-error {{same pointer dereferenced in multiple different ways in map clause expressions}} lt50-note {{used here}}
{}
- #pragma omp target teams map(r.RPtrS, r.RPtrS->B) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
+ #pragma omp target teams map(r.RPtrS, r.RPtrS->B) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}}
{}
#pragma omp target teams map(r.RPtrS->A, r.RPtrS->B)
{}
@@ -262,13 +262,13 @@
{}
#pragma omp target teams map(r.C, r.D)
{}
- #pragma omp target teams map(r.C, r.C) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+ #pragma omp target teams map(r.C, r.C) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
{}
- #pragma omp target teams map(r.C) map(r.C) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+ #pragma omp target teams map(r.C) map(r.C) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
{}
#pragma omp target teams map(r.C, r.S) // this would be an error only caught at runtime - Sema would have to make sure there is not way for the missing data between fields to be mapped somewhere else.
{}
- #pragma omp target teams map(r, r.S) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+ #pragma omp target teams map(r, r.S) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
{}
#pragma omp target teams map(r.C, t.C)
{}
@@ -298,15 +298,15 @@
#pragma omp target teams map(u.B) // expected-error {{mapping of union members is not allowed}}
{}
- #pragma omp target data map(to: r.C) //expected-note {{used here}}
+ #pragma omp target data map(to: r.C) // lt50-note {{used here}}
{
- #pragma omp target teams map(r.D) // expected-error {{original storage of expression in data environment is shared but data environment do not fully contain mapped expression storage}}
+ #pragma omp target teams map(r.D) // lt50-error {{original storage of expression in data environment is shared but data environment do not fully contain mapped expression storage}}
{}
}
- #pragma omp target data map(to: t.Ptr) //expected-note {{used here}}
+ #pragma omp target data map(to: t.Ptr) // lt50-note {{used here}}
{
- #pragma omp target teams map(t.Ptr[:23]) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+ #pragma omp target teams map(t.Ptr[:23]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
{}
}
@@ -457,22 +457,22 @@
#pragma omp target data map(S2::S2sc)
#pragma omp target data map(e, g)
#pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
-#pragma omp target data map(k) map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
-#pragma omp target teams map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}}
+#pragma omp target data map(k) map(k) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
+#pragma omp target teams map(k), map(k[:5]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}} lt50-note 2 {{used here}}
foo();
#pragma omp target data map(da)
#pragma omp target teams map(da[:4])
foo();
-#pragma omp target data map(k, j, l) // expected-note 2 {{used here}}
-#pragma omp target data map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l) // lt50-note 2 {{used here}}
+#pragma omp target data map(k[:4]) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
#pragma omp target data map(j)
-#pragma omp target teams map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+#pragma omp target teams map(l) map(l[:5]) // lt50-error 2 {{variable already marked as mapped in current construct}} lt50-note 2 {{used here}}
foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}}
-#pragma omp target data map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note 2 {{used here}}
+#pragma omp target data map(k) // lt50-error 2 {{pointer cannot be mapped along with a section derived from itself}}
#pragma omp target data map(j)
#pragma omp target teams map(l)
foo();
@@ -540,22 +540,22 @@
#pragma omp target data map(S2::S2sc)
#pragma omp target data map(e, g)
#pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
-#pragma omp target data map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
-#pragma omp target teams map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
+#pragma omp target data map(k), map(k) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
+#pragma omp target teams map(k), map(k[:5]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}} lt50-note {{used here}}
foo();
#pragma omp target data map(da)
#pragma omp target teams map(da[:4])
foo();
-#pragma omp target data map(k, j, l) // expected-note {{used here}}
-#pragma omp target data map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k, j, l) // lt50-note {{used here}}
+#pragma omp target data map(k[:4]) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
#pragma omp target data map(j)
-#pragma omp target teams map(l) map(l[:5]) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+#pragma omp target teams map(l) map(l[:5]) // lt50-error {{variable already marked as mapped in current construct}} lt50-note {{used here}}
foo();
-#pragma omp target data map(k[:4], j, l[:5]) // expected-note {{used here}}
-#pragma omp target data map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(k[:4], j, l[:5]) // lt50-note {{used here}}
+#pragma omp target data map(k) // lt50-error {{pointer cannot be mapped along with a section derived from itself}}
#pragma omp target data map(j)
#pragma omp target teams map(l)
foo();
diff --git a/clang/test/OpenMP/target_update_codegen.cpp b/clang/test/OpenMP/target_update_codegen.cpp
--- a/clang/test/OpenMP/target_update_codegen.cpp
+++ b/clang/test/OpenMP/target_update_codegen.cpp
@@ -30,15 +30,15 @@
double gc[100];
// CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 800]
-// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 34]
+// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 2]
// CK1: [[SIZE02:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 4]
-// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 1]
-// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 34]
+// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 2]
// CK1: [[SIZE04:@.+]] = {{.+}}constant [2 x i64] [i64 sdiv exact (i64 sub (i64 ptrtoint (double** getelementptr (double*, double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1), i32 1) to i64), i64 ptrtoint (double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1) to i64)), i64 ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)), i64 24]
-// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673]
+// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710673]
// CK1-LABEL: _Z3fooi
void foo(int arg) {
@@ -174,7 +174,7 @@
}
};
-// CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710674]
+// CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710674]
// CK2-LABEL: _Z3bari
int bar(int arg){
@@ -311,7 +311,7 @@
#ifdef CK5
// CK5: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
-// CK5: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK5: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 1]
// CK5-LABEL: lvalue
void lvalue(int *B, int l, int e) {
@@ -353,7 +353,7 @@
#ifdef CK6
// CK6: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
-// CK6: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK6: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 1]
// CK6-LABEL: lvalue
void lvalue(int *B, int l, int e) {
@@ -400,7 +400,7 @@
#ifdef CK7
// CK7: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
-// CK7: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK7: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 1]
// CK7-LABEL: lvalue
void lvalue(int *B, int l, int e) {
@@ -450,7 +450,7 @@
#ifdef CK8
// CK8: [[SIZE00:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}, i{{64|32}} 4]
-// CK8: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 33, i64 17]
+// CK8: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 1, i64 17]
// CK8-LABEL: lvalue
void lvalue(int **B, int l, int e) {
@@ -504,7 +504,7 @@
double *p;
};
-// CK9: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673]
+// CK9: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710673]
// CK9-LABEL: lvalue
void lvalue(struct S *s, int l, int e) {
@@ -554,7 +554,7 @@
double *p;
};
-// CK10: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673]
+// CK10: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710673]
// CK10-LABEL: lvalue
void lvalue(struct S *s, int l, int e) {
@@ -604,7 +604,7 @@
struct S {
double *p;
};
-// CK11: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673]
+// CK11: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710673]
// CK11-LABEL: lvalue
void lvalue(struct S *s, int l, int e) {
@@ -656,7 +656,7 @@
double *p;
struct S *sp;
};
-// CK12: [[MTYPE00:@.+]] = {{.+}}constant [3 x i64] [i64 32, i64 281474976710672, i64 17]
+// CK12: [[MTYPE00:@.+]] = {{.+}}constant [3 x i64] [i64 0, i64 281474976710672, i64 17]
// CK12-LABEL: lvalue
void lvalue(struct S *s, int l, int e) {
@@ -718,7 +718,7 @@
#ifdef CK13
// CK13: [[SIZE00:@.+]] = {{.+}}constant [1 x i64] [i64 4]
-// CK13: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK13: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 1]
// CK13-LABEL: lvalue
void lvalue(int **BB, int a, int b) {
@@ -765,7 +765,7 @@
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
#ifdef CK14
-// CK14: [[MTYPE00:@.+]] = private {{.*}}constant [2 x i64] [i64 32, i64 281474976710673]
+// CK14: [[MTYPE00:@.+]] = private {{.*}}constant [2 x i64] [i64 0, i64 281474976710673]
struct SSA {
double *p;
@@ -838,7 +838,7 @@
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
#ifdef CK15
-// CK15: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673]
+// CK15: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710673]
struct SSA {
double *p;
@@ -904,7 +904,7 @@
#ifdef CK16
// CK16: [[SIZE00:@.+]] = {{.+}}constant [1 x i64] [i64 4]
-// CK16: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK16: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 1]
//CK16-LABEL: lvalue_find_base
void lvalue_find_base(float *f, int *i) {
@@ -949,7 +949,7 @@
#ifdef CK17
// CK17: [[SIZE00:@.+]] = {{.+}}constant [1 x i64] [i64 4]
-// CK17: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK17: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 1]
struct SSA {
int i;
@@ -1007,8 +1007,8 @@
// SIMD-ONLY18-NOT: {{__kmpc|__tgt}}
#ifdef CK18
-// CK18-DAG: [[MTYPE_TO:@.+]] = {{.+}}constant [1 x i64] [i64 33]
-// CK18-DAG: [[MTYPE_FROM:@.+]] = {{.+}}constant [1 x i64] [i64 34]
+// CK18-DAG: [[MTYPE_TO:@.+]] = {{.+}}constant [1 x i64] [i64 1]
+// CK18-DAG: [[MTYPE_FROM:@.+]] = {{.+}}constant [1 x i64] [i64 2]
//CK18-LABEL: array_shaping
void array_shaping(float *f, int sa) {
@@ -1082,11 +1082,11 @@
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
#ifdef CK19
-// PRESENT=0x1000 | TARGET_PARAM=0x20 | TO=0x1 = 0x1021
-// CK19: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1021]]]
+// PRESENT=0x1000 | TO=0x1 = 0x1001
+// CK19: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1001]]]
-// PRESENT=0x1000 | TARGET_PARAM=0x20 | FROM=0x2 = 0x1022
-// CK19: [[MTYPE01:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1022]]]
+// PRESENT=0x1000 | FROM=0x2 = 0x1022
+// CK19: [[MTYPE01:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1002]]]
// CK19-LABEL: _Z13check_presenti
void check_present(int arg) {
diff --git a/clang/test/OpenMP/target_update_depend_codegen.cpp b/clang/test/OpenMP/target_update_depend_codegen.cpp
--- a/clang/test/OpenMP/target_update_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_update_depend_codegen.cpp
@@ -30,15 +30,15 @@
double gc[100];
// CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i64] [i64 800]
-// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 34]
+// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 2]
// CK1: [[SIZE02:@.+]] = {{.+}}constant [1 x i64] [i64 4]
-// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 1]
-// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 34]
+// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 2]
// CK1: [[SIZE04:@.+]] = {{.+}}constant [2 x i64] [i64 sdiv exact (i64 sub (i64 ptrtoint (double** getelementptr (double*, double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1), i32 1) to i64), i64 ptrtoint (double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1) to i64)), i64 ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)), i64 24]
-// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673]
+// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710673]
// CK1-LABEL: _Z3fooi
void foo(int arg) {
diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp
--- a/openmp/libomptarget/src/omptarget.cpp
+++ b/openmp/libomptarget/src/omptarget.cpp
@@ -230,7 +230,7 @@
for (unsigned I = 0, E = MapperComponents.Components.size(); I < E; ++I) {
auto &C =
MapperComponents
- .Components[target_data_function == targetDataEnd ? I : E - I - 1];
+ .Components[target_data_function == targetDataEnd ? E - I - 1 : I];
MapperArgsBase[I] = C.Base;
MapperArgs[I] = C.Begin;
MapperArgSizes[I] = C.Size;