[clang] 0caf736 - [OPENMP50]Mapping of the subcomponents with the 'default' mappers.
Alexey Bataev via cfe-commits
cfe-commits at lists.llvm.org
Tue Mar 2 07:14:34 PST 2021
Author: Alexey Bataev
Date: 2021-03-02T07:11:06-08:00
New Revision: 0caf736d7e1d16d1059553fc28dbac31f0b9f788
URL: https://github.com/llvm/llvm-project/commit/0caf736d7e1d16d1059553fc28dbac31f0b9f788
DIFF: https://github.com/llvm/llvm-project/commit/0caf736d7e1d16d1059553fc28dbac31f0b9f788.diff
LOG: [OPENMP50]Mapping of the subcomponents with the 'default' mappers.
If the mapped structure has data members, which have 'default' mappers,
need to map these members individually using their 'default' mappers.
Differential Revision: https://reviews.llvm.org/D92195
Added:
clang/test/OpenMP/target_map_codegen_34.cpp
openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers.cpp
Modified:
clang/include/clang/AST/OpenMPClause.h
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/lib/Sema/SemaOpenMP.cpp
clang/test/OpenMP/declare_mapper_codegen.cpp
openmp/libomptarget/src/omptarget.cpp
openmp/libomptarget/src/private.h
Removed:
################################################################################
diff --git a/clang/include/clang/AST/OpenMPClause.h b/clang/include/clang/AST/OpenMPClause.h
index 9c6bbe6082fb..958f2b9e0e6f 100644
--- a/clang/include/clang/AST/OpenMPClause.h
+++ b/clang/include/clang/AST/OpenMPClause.h
@@ -5355,14 +5355,14 @@ class OMPMappableExprListClause : public OMPVarListClause<T>,
if (!(--RemainingLists)) {
++DeclCur;
++NumListsCur;
- if (SupportsMapper)
- ++MapperCur;
RemainingLists = *NumListsCur;
assert(RemainingLists && "No lists in the following declaration??");
}
}
++ListSizeCur;
+ if (SupportsMapper)
+ ++MapperCur;
return *this;
}
};
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index a56dad819ac0..d0876056268d 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7151,11 +7151,13 @@ class MappableExprsHandler {
/// [ValueDecl *] --> {LE(FieldIndex, Pointer),
/// HE(FieldIndex, Pointer)}
struct StructRangeInfoTy {
+ MapCombinedInfoTy PreliminaryMapData;
std::pair<unsigned /*FieldIndex*/, Address /*Pointer*/> LowestElem = {
0, Address::invalid()};
std::pair<unsigned /*FieldIndex*/, Address /*Pointer*/> HighestElem = {
0, Address::invalid()};
Address Base = Address::invalid();
+ Address LB = Address::invalid();
bool IsArraySection = false;
bool HasCompleteRecord = false;
};
@@ -7754,11 +7756,9 @@ class MappableExprsHandler {
(IsPointer || ForDeviceAddr) && EncounteredME &&
(dyn_cast<MemberExpr>(I->getAssociatedExpression()) ==
EncounteredME);
- if (!OverlappedElements.empty()) {
+ if (!OverlappedElements.empty() && Next == CE) {
// Handle base element with the info for overlapped elements.
assert(!PartialStruct.Base.isValid() && "The base element is set.");
- assert(Next == CE &&
- "Expected last element for the overlapped elements.");
assert(!IsPointer &&
"Unexpected base element with the pointer type.");
// Mark the whole struct as the struct that requires allocation on the
@@ -7775,13 +7775,17 @@ class MappableExprsHandler {
PartialStruct.HighestElem.first)>::max(),
HB};
PartialStruct.Base = BP;
+ PartialStruct.LB = LB;
+ assert(
+ PartialStruct.PreliminaryMapData.BasePointers.empty() &&
+ "Overlapped elements must be used only once for the variable.");
+ std::swap(PartialStruct.PreliminaryMapData, CombinedInfo);
// Emit data for non-overlapped data.
OpenMPOffloadMappingFlags Flags =
OMP_MAP_MEMBER_OF |
getMapTypeBits(MapType, MapModifiers, MotionModifiers, IsImplicit,
/*AddPtrFlag=*/false,
/*AddIsTargetParamFlag=*/false, IsNonContiguous);
- LB = BP;
llvm::Value *Size = nullptr;
// Do bitcopy of all non-overlapped structure elements.
for (OMPClauseMappableExprCommon::MappableExprComponentListRef
@@ -7890,6 +7894,7 @@ class MappableExprsHandler {
PartialStruct.HighestElem = {FieldIndex, LB};
}
PartialStruct.Base = BP;
+ PartialStruct.LB = BP;
} else if (FieldIndex < PartialStruct.LowestElem.first) {
PartialStruct.LowestElem = {FieldIndex, LB};
} else if (FieldIndex > PartialStruct.HighestElem.first) {
@@ -8609,8 +8614,8 @@ class MappableExprsHandler {
Address LBAddr = PartialStruct.LowestElem.second;
Address HBAddr = PartialStruct.HighestElem.second;
if (PartialStruct.HasCompleteRecord) {
- LBAddr = PartialStruct.Base;
- HBAddr = PartialStruct.Base;
+ LBAddr = PartialStruct.LB;
+ HBAddr = PartialStruct.LB;
}
CombinedInfo.Exprs.push_back(VD);
// Base is the base of the struct
@@ -8909,11 +8914,17 @@ class MappableExprsHandler {
// Sort the overlapped elements for each item.
llvm::SmallVector<const FieldDecl *, 4> Layout;
if (!OverlappedData.empty()) {
- if (const auto *CRD =
- VD->getType().getCanonicalType()->getAsCXXRecordDecl())
+ const Type *BaseType = VD->getType().getCanonicalType().getTypePtr();
+ const Type *OrigType = BaseType->getPointeeOrArrayElementType();
+ while (BaseType != OrigType) {
+ BaseType = OrigType->getCanonicalTypeInternal().getTypePtr();
+ OrigType = BaseType->getPointeeOrArrayElementType();
+ }
+
+ if (const auto *CRD = BaseType->getAsCXXRecordDecl())
getPlainLayout(CRD, Layout, /*AsBase=*/false);
else {
- const auto *RD = VD->getType().getCanonicalType()->getAsRecordDecl();
+ const auto *RD = BaseType->getAsRecordDecl();
Layout.append(RD->field_begin(), RD->field_end());
}
}
@@ -9567,10 +9578,12 @@ getNestedDistributeDirective(ASTContext &Ctx, const OMPExecutableDirective &D) {
/// void *base, void *begin,
/// int64_t size, int64_t type,
/// void *name = nullptr) {
-/// // Allocate space for an array section first.
-/// if ((size > 1 || base != begin) && !maptype.IsDelete)
+/// // Allocate space for an array section first or add a base/begin for
+/// // pointer dereference.
+/// if ((size > 1 || (base != begin && maptype.IsPtrAndObj)) &&
+/// !maptype.IsDelete)
/// __tgt_push_mapper_component(rt_mapper_handle, base, begin,
-/// size*sizeof(Ty), clearToFrom(type));
+/// size*sizeof(Ty), clearToFromMember(type));
/// // Map members.
/// for (unsigned i = 0; i < size; i++) {
/// // For each component specified by this mapper:
@@ -9585,9 +9598,9 @@ getNestedDistributeDirective(ASTContext &Ctx, const OMPExecutableDirective &D) {
/// }
/// }
/// // Delete the array section.
-/// if ((size > 1 || base != begin) && maptype.IsDelete)
+/// if (size > 1 && maptype.IsDelete)
/// __tgt_push_mapper_component(rt_mapper_handle, base, begin,
-/// size*sizeof(Ty), clearToFrom(type));
+/// size*sizeof(Ty), clearToFromMember(type));
/// }
/// \endcode
void CGOpenMPRuntime::emitUserDefinedMapper(const OMPDeclareMapperDecl *D,
@@ -9851,18 +9864,26 @@ void CGOpenMPRuntime::emitUDMapperArrayInitOrDel(
MapperCGF.createBasicBlock(getName({"omp.array", Prefix}));
llvm::Value *IsArray = MapperCGF.Builder.CreateICmpSGT(
Size, MapperCGF.Builder.getInt64(1), "omp.arrayinit.isarray");
- // base != begin?
- llvm::Value *BaseIsBegin = MapperCGF.Builder.CreateIsNotNull(
- MapperCGF.Builder.CreatePtrDiff(Base, Begin));
- llvm::Value *Cond = MapperCGF.Builder.CreateOr(IsArray, BaseIsBegin);
llvm::Value *DeleteBit = MapperCGF.Builder.CreateAnd(
MapType,
MapperCGF.Builder.getInt64(MappableExprsHandler::OMP_MAP_DELETE));
llvm::Value *DeleteCond;
+ llvm::Value *Cond;
if (IsInit) {
+ // base != begin?
+ llvm::Value *BaseIsBegin = MapperCGF.Builder.CreateIsNotNull(
+ MapperCGF.Builder.CreatePtrDiff(Base, Begin));
+ // IsPtrAndObj?
+ llvm::Value *PtrAndObjBit = MapperCGF.Builder.CreateAnd(
+ MapType,
+ MapperCGF.Builder.getInt64(MappableExprsHandler::OMP_MAP_PTR_AND_OBJ));
+ PtrAndObjBit = MapperCGF.Builder.CreateIsNotNull(PtrAndObjBit);
+ BaseIsBegin = MapperCGF.Builder.CreateAnd(BaseIsBegin, PtrAndObjBit);
+ Cond = MapperCGF.Builder.CreateOr(IsArray, BaseIsBegin);
DeleteCond = MapperCGF.Builder.CreateIsNull(
DeleteBit, getName({"omp.array", Prefix, ".delete"}));
} else {
+ Cond = IsArray;
DeleteCond = MapperCGF.Builder.CreateIsNotNull(
DeleteBit, getName({"omp.array", Prefix, ".delete"}));
}
@@ -9879,7 +9900,8 @@ void CGOpenMPRuntime::emitUDMapperArrayInitOrDel(
llvm::Value *MapTypeArg = MapperCGF.Builder.CreateAnd(
MapType,
MapperCGF.Builder.getInt64(~(MappableExprsHandler::OMP_MAP_TO |
- MappableExprsHandler::OMP_MAP_FROM)));
+ MappableExprsHandler::OMP_MAP_FROM |
+ MappableExprsHandler::OMP_MAP_MEMBER_OF)));
llvm::Value *MapNameArg = llvm::ConstantPointerNull::get(CGM.VoidPtrTy);
// Call the runtime API __tgt_push_mapper_component to fill up the runtime
@@ -10171,9 +10193,12 @@ void CGOpenMPRuntime::emitTargetCall(
// 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,
- nullptr, /*NoTargetParam=*/false);
+ if (PartialStruct.Base.isValid()) {
+ CombinedInfo.append(PartialStruct.PreliminaryMapData);
+ MEHandler.emitCombinedEntry(
+ CombinedInfo, CurInfo.Types, PartialStruct, nullptr,
+ !PartialStruct.PreliminaryMapData.BasePointers.empty());
+ }
// We need to append the results of this capture to what we already have.
CombinedInfo.append(CurInfo);
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index f2c9fe644d96..0bb9554bcab8 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -5151,6 +5151,146 @@ static void checkAllocateClauses(Sema &S, DSAStackTy *Stack,
}
}
+static ExprResult buildUserDefinedMapperRef(Sema &SemaRef, Scope *S,
+ CXXScopeSpec &MapperIdScopeSpec,
+ const DeclarationNameInfo &MapperId,
+ QualType Type,
+ Expr *UnresolvedMapper);
+
+/// Perform DFS through the structure/class data members trying to find
+/// member(s) with user-defined 'default' mapper and generate implicit map
+/// clauses for such members with the found 'default' mapper.
+static void
+processImplicitMapsWithDefaultMappers(Sema &S, DSAStackTy *Stack,
+ SmallVectorImpl<OMPClause *> &Clauses) {
+ // Check for the deault mapper for data members.
+ if (S.getLangOpts().OpenMP < 50)
+ return;
+ SmallVector<OMPClause *, 4> ImplicitMaps;
+ DeclarationNameInfo DefaultMapperId;
+ DefaultMapperId.setName(S.Context.DeclarationNames.getIdentifier(
+ &S.Context.Idents.get("default")));
+ for (int Cnt = 0, EndCnt = Clauses.size(); Cnt < EndCnt; ++Cnt) {
+ auto *C = dyn_cast<OMPMapClause>(Clauses[Cnt]);
+ if (!C)
+ continue;
+ SmallVector<Expr *, 4> SubExprs;
+ auto *MI = C->mapperlist_begin();
+ for (auto I = C->varlist_begin(), End = C->varlist_end(); I != End;
+ ++I, ++MI) {
+ // Expression is mapped using mapper - skip it.
+ if (*MI)
+ continue;
+ Expr *E = *I;
+ // Expression is dependent - skip it, build the mapper when it gets
+ // instantiated.
+ if (E->isTypeDependent() || E->isValueDependent() ||
+ E->containsUnexpandedParameterPack())
+ continue;
+ // Array section - need to check for the mapping of the array section
+ // element.
+ QualType CanonType = E->getType().getCanonicalType();
+ if (CanonType->isSpecificBuiltinType(BuiltinType::OMPArraySection)) {
+ const auto *OASE = cast<OMPArraySectionExpr>(E->IgnoreParenImpCasts());
+ QualType BaseType =
+ OMPArraySectionExpr::getBaseOriginalType(OASE->getBase());
+ QualType ElemType;
+ if (const auto *ATy = BaseType->getAsArrayTypeUnsafe())
+ ElemType = ATy->getElementType();
+ else
+ ElemType = BaseType->getPointeeType();
+ CanonType = ElemType;
+ }
+
+ // DFS over data members in structures/classes.
+ SmallVector<std::pair<QualType, FieldDecl *>, 4> Types(
+ 1, {CanonType, nullptr});
+ llvm::DenseMap<const Type *, Expr *> Visited;
+ SmallVector<std::pair<FieldDecl *, unsigned>, 4> ParentChain(
+ 1, {nullptr, 1});
+ while (!Types.empty()) {
+ QualType BaseType;
+ FieldDecl *CurFD;
+ std::tie(BaseType, CurFD) = Types.pop_back_val();
+ while (ParentChain.back().second == 0)
+ ParentChain.pop_back();
+ --ParentChain.back().second;
+ if (BaseType.isNull())
+ continue;
+ // Only structs/classes are allowed to have mappers.
+ const RecordDecl *RD = BaseType.getCanonicalType()->getAsRecordDecl();
+ if (!RD)
+ continue;
+ auto It = Visited.find(BaseType.getTypePtr());
+ if (It == Visited.end()) {
+ // Try to find the associated user-defined mapper.
+ CXXScopeSpec MapperIdScopeSpec;
+ ExprResult ER = buildUserDefinedMapperRef(
+ S, Stack->getCurScope(), MapperIdScopeSpec, DefaultMapperId,
+ BaseType, /*UnresolvedMapper=*/nullptr);
+ if (ER.isInvalid())
+ continue;
+ It = Visited.try_emplace(BaseType.getTypePtr(), ER.get()).first;
+ }
+ // Found default mapper.
+ if (It->second) {
+ auto *OE = new (S.Context) OpaqueValueExpr(E->getExprLoc(), CanonType,
+ VK_LValue, OK_Ordinary, E);
+ OE->setIsUnique(/*V=*/true);
+ Expr *BaseExpr = OE;
+ for (const auto &P : ParentChain) {
+ if (P.first) {
+ BaseExpr = S.BuildMemberExpr(
+ BaseExpr, /*IsArrow=*/false, E->getExprLoc(),
+ NestedNameSpecifierLoc(), SourceLocation(), P.first,
+ DeclAccessPair::make(P.first, P.first->getAccess()),
+ /*HadMultipleCandidates=*/false, DeclarationNameInfo(),
+ P.first->getType(), VK_LValue, OK_Ordinary);
+ BaseExpr = S.DefaultLvalueConversion(BaseExpr).get();
+ }
+ }
+ if (CurFD)
+ BaseExpr = S.BuildMemberExpr(
+ BaseExpr, /*IsArrow=*/false, E->getExprLoc(),
+ NestedNameSpecifierLoc(), SourceLocation(), CurFD,
+ DeclAccessPair::make(CurFD, CurFD->getAccess()),
+ /*HadMultipleCandidates=*/false, DeclarationNameInfo(),
+ CurFD->getType(), VK_LValue, OK_Ordinary);
+ SubExprs.push_back(BaseExpr);
+ continue;
+ }
+ // Check for the "default" mapper for data memebers.
+ bool FirstIter = true;
+ for (FieldDecl *FD : RD->fields()) {
+ if (!FD)
+ continue;
+ QualType FieldTy = FD->getType();
+ if (FieldTy.isNull() ||
+ !(FieldTy->isStructureOrClassType() || FieldTy->isUnionType()))
+ continue;
+ if (FirstIter) {
+ FirstIter = false;
+ ParentChain.emplace_back(CurFD, 1);
+ } else {
+ ++ParentChain.back().second;
+ }
+ Types.emplace_back(FieldTy, FD);
+ }
+ }
+ }
+ if (SubExprs.empty())
+ continue;
+ CXXScopeSpec MapperIdScopeSpec;
+ DeclarationNameInfo MapperId;
+ if (OMPClause *NewClause = S.ActOnOpenMPMapClause(
+ C->getMapTypeModifiers(), C->getMapTypeModifiersLoc(),
+ MapperIdScopeSpec, MapperId, C->getMapType(),
+ /*IsMapTypeImplicit=*/true, SourceLocation(), SourceLocation(),
+ SubExprs, OMPVarListLocTy()))
+ Clauses.push_back(NewClause);
+ }
+}
+
StmtResult Sema::ActOnOpenMPExecutableDirective(
OpenMPDirectiveKind Kind, const DeclarationNameInfo &DirName,
OpenMPDirectiveKind CancelRegion, ArrayRef<OMPClause *> Clauses,
@@ -5271,6 +5411,11 @@ StmtResult Sema::ActOnOpenMPExecutableDirective(
}
}
}
+ // Build expressions for implicit maps of data members with 'default'
+ // mappers.
+ if (LangOpts.OpenMP >= 50)
+ processImplicitMapsWithDefaultMappers(*this, DSAStack,
+ ClausesWithImplicit);
}
llvm::SmallVector<OpenMPDirectiveKind, 4> AllowedNameModifiers;
@@ -17502,6 +17647,14 @@ class MapBaseChecker final : public StmtVisitor<MapBaseChecker, bool> {
Components.emplace_back(COCE, nullptr, IsNonContiguous);
return true;
}
+ bool VisitOpaqueValueExpr(OpaqueValueExpr *E) {
+ Expr *Source = E->getSourceExpr();
+ if (!Source) {
+ emitErrorMsg();
+ return false;
+ }
+ return Visit(Source);
+ }
bool VisitStmt(Stmt *) {
emitErrorMsg();
return false;
@@ -18622,8 +18775,15 @@ Sema::DeclGroupPtrTy Sema::ActOnOpenMPDeclareMapperDirective(
Diag(I->second, diag::note_previous_definition);
Invalid = true;
}
- auto *DMD = OMPDeclareMapperDecl::Create(Context, DC, StartLoc, Name,
- MapperType, VN, Clauses, PrevDMD);
+ // Build expressions for implicit maps of data members with 'default'
+ // mappers.
+ SmallVector<OMPClause *, 4> ClausesWithImplicit(Clauses.begin(),
+ Clauses.end());
+ if (LangOpts.OpenMP >= 50)
+ processImplicitMapsWithDefaultMappers(*this, DSAStack, ClausesWithImplicit);
+ auto *DMD =
+ OMPDeclareMapperDecl::Create(Context, DC, StartLoc, Name, MapperType, VN,
+ ClausesWithImplicit, PrevDMD);
if (S)
PushOnScopeChains(DMD, S);
else
diff --git a/clang/test/OpenMP/declare_mapper_codegen.cpp b/clang/test/OpenMP/declare_mapper_codegen.cpp
index 8d820e8e3355..7ccb361a2d0a 100644
--- a/clang/test/OpenMP/declare_mapper_codegen.cpp
+++ b/clang/test/OpenMP/declare_mapper_codegen.cpp
@@ -107,7 +107,10 @@ class C {
// CK0-DAG: [[DIF:%.+]] = sub i64 [[BPTRI]], [[PTRI]]
// CK0-DAG: [[NORM:%.+]] = sdiv exact i64 [[DIF]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
// CK0-DAG: [[PTRSNE:%.+]] = icmp ne i64 [[NORM]], 0
-// CK0-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[PTRSNE]]
+// CK0-DAG: [[PTRANDOBJ:%.+]] = and i64 [[TYPE]], 16
+// CK0-DAG: [[ISPTRANDOBJ:%.+]] = icmp ne i64 [[PTRANDOBJ]], 0
+// CK0-DAG: [[CMPA:%.+]] = and i1 [[PTRSNE]], [[ISPTRANDOBJ]]
+// CK0-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[CMPA]]
// CK0-DAG: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
// CK0-DAG: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0
// CK0-DAG: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]]
@@ -115,7 +118,7 @@ class C {
// 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-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], 281474976710652
// CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]], {{.*}})
// CK0: br label %[[LHEAD:[^,]+]]
@@ -218,20 +221,14 @@ class C {
// CK0: [[LEXIT]]
// CK0: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1
-// CK0: [[BPTRI:%.+]] = ptrtoint i8* [[BPTR]] to i64
-// CK0: [[PTRI:%.+]] = ptrtoint i8* [[BEGIN]] to i64
-// CK0: [[DIF:%.+]] = sub i64 [[BPTRI]], [[PTRI]]
-// CK0: [[NORM:%.+]] = sdiv exact i64 [[DIF]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
-// CK0: [[PTRSNE:%.+]] = icmp ne i64 [[NORM]], 0
-// CK0: [[CMP:%.+]] = or i1 [[ISARRAY]], [[PTRSNE]]
// CK0: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
// CK0: [[ISNOTDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
-// CK0: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]]
+// CK0: [[CMP1:%.+]] = and i1 [[ISARRAY]], [[ISNOTDEL]]
// CK0: br i1 [[CMP1]], label %[[EVALDEL:[^,]+]], label %[[DONE]]
// CK0: [[EVALDEL]]
// CK0-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
// CK0-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
-// CK0-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
+// CK0-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], 281474976710652
// CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]], {{.*}})
// CK0: br label %[[DONE]]
// CK0: [[DONE]]
@@ -659,7 +656,10 @@ class C {
// CK1-DAG: [[DIF:%.+]] = sub i64 [[BPTRI]], [[PTRI]]
// CK1-DAG: [[NORM:%.+]] = sdiv exact i64 [[DIF]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
// CK1-DAG: [[PTRSNE:%.+]] = icmp ne i64 [[NORM]], 0
-// CK1-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[PTRSNE]]
+// CK1-DAG: [[PTRANDOBJ:%.+]] = and i64 [[TYPE]], 16
+// CK1-DAG: [[ISPTRANDOBJ:%.+]] = icmp ne i64 [[PTRANDOBJ]], 0
+// CK1-DAG: [[CMPA:%.+]] = and i1 [[PTRSNE]], [[ISPTRANDOBJ]]
+// CK1-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[CMPA]]
// CK1-DAG: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
// CK1-DAG: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0
// CK1-DAG: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]]
@@ -667,7 +667,7 @@ class C {
// CK1: [[INITEVALDEL]]
// CK1-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 4
-// CK1-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
+// CK1-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], 281474976710652
// CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]], {{.*}})
// CK1: br label %[[LHEAD:[^,]+]]
@@ -709,17 +709,11 @@ class C {
// CK1: [[LEXIT]]
// CK1: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1
-// CK1: [[BPTRI:%.+]] = ptrtoint i8* [[BPTR]] to i64
-// CK1: [[PTRI:%.+]] = ptrtoint i8* [[BEGIN]] to i64
-// CK1: [[DIF:%.+]] = sub i64 [[BPTRI]], [[PTRI]]
-// CK1: [[NORM:%.+]] = sdiv exact i64 [[DIF]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
-// CK1: [[PTRSNE:%.+]] = icmp ne i64 [[NORM]], 0
-// CK1: [[CMP:%.+]] = or i1 [[ISARRAY]], [[PTRSNE]]
// CK1: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
// CK1: [[ISNOTDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
-// CK1: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]]
+// CK1: [[CMP1:%.+]] = and i1 [[ISARRAY]], [[ISNOTDEL]]
// CK1-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 4
-// CK1-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
+// CK1-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], 281474976710652
// CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]], {{.*}})
// CK1: br label %[[DONE]]
// CK1: [[DONE]]
@@ -783,7 +777,10 @@ class C {
// CK2-DAG: [[DIF:%.+]] = sub i64 [[BPTRI]], [[PTRI]]
// CK2-DAG: [[NORM:%.+]] = sdiv exact i64 [[DIF]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
// CK2-DAG: [[PTRSNE:%.+]] = icmp ne i64 [[NORM]], 0
-// CK2-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[PTRSNE]]
+// CK2-DAG: [[PTRANDOBJ:%.+]] = and i64 [[TYPE]], 16
+// CK2-DAG: [[ISPTRANDOBJ:%.+]] = icmp ne i64 [[PTRANDOBJ]], 0
+// CK2-DAG: [[CMPA:%.+]] = and i1 [[PTRSNE]], [[ISPTRANDOBJ]]
+// CK2-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[CMPA]]
// CK2-DAG: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
// CK2-DAG: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0
// CK2-DAG: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]]
@@ -791,7 +788,7 @@ class C {
// CK2: [[INITEVALDEL]]
// CK2-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
-// CK2-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
+// CK2-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], 281474976710652
// CK2: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]], {{.*}})
// CK2: br label %[[LHEAD:[^,]+]]
@@ -833,19 +830,13 @@ class C {
// CK2: [[LEXIT]]
// CK2: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1
-// CK2: [[BPTRI:%.+]] = ptrtoint i8* [[BPTR]] to i64
-// CK2: [[PTRI:%.+]] = ptrtoint i8* [[BEGIN]] to i64
-// CK2: [[DIF:%.+]] = sub i64 [[BPTRI]], [[PTRI]]
-// CK2: [[NORM:%.+]] = sdiv exact i64 [[DIF]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
-// CK2: [[PTRSNE:%.+]] = icmp ne i64 [[NORM]], 0
-// CK2: [[CMP:%.+]] = or i1 [[ISARRAY]], [[PTRSNE]]
// CK2: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
// CK2: [[ISNOTDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
-// CK2: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]]
+// CK2: [[CMP1:%.+]] = and i1 [[ISARRAY]], [[ISNOTDEL]]
// CK2: br i1 [[CMP1]], label %[[EVALDEL:[^,]+]], label %[[DONE]]
// CK2: [[EVALDEL]]
// CK2-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
-// CK2-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
+// CK2-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], 281474976710652
// CK2: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]], {{.*}})
// CK2: br label %[[DONE]]
// CK2: [[DONE]]
@@ -990,7 +981,10 @@ class C {
// CK4-DAG: [[DIF:%.+]] = sub i64 [[BPTRI]], [[PTRI]]
// CK4-DAG: [[NORM:%.+]] = sdiv exact i64 [[DIF]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
// CK4-DAG: [[PTRSNE:%.+]] = icmp ne i64 [[NORM]], 0
-// CK4-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[PTRSNE]]
+// CK4-DAG: [[PTRANDOBJ:%.+]] = and i64 [[TYPE]], 16
+// CK4-DAG: [[ISPTRANDOBJ:%.+]] = icmp ne i64 [[PTRANDOBJ]], 0
+// CK4-DAG: [[CMPA:%.+]] = and i1 [[PTRSNE]], [[ISPTRANDOBJ]]
+// CK4-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[CMPA]]
// CK4-DAG: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
// CK4-DAG: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0
// CK4-DAG: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]]
@@ -999,7 +993,7 @@ class C {
// CK4: [[INITEVALDEL]]
// 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-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], 281474976710652
// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]], {{.*}})
// CK4: br label %[[LHEAD:[^,]+]]
@@ -1102,20 +1096,14 @@ class C {
// CK4: [[LEXIT]]
// CK4: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1
-// CK4: [[BPTRI:%.+]] = ptrtoint i8* [[BPTR]] to i64
-// CK4: [[PTRI:%.+]] = ptrtoint i8* [[BEGIN]] to i64
-// CK4: [[DIF:%.+]] = sub i64 [[BPTRI]], [[PTRI]]
-// CK4: [[NORM:%.+]] = sdiv exact i64 [[DIF]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
-// CK4: [[PTRSNE:%.+]] = icmp ne i64 [[NORM]], 0
-// CK4: [[CMP:%.+]] = or i1 [[ISARRAY]], [[PTRSNE]]
// CK4: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
// CK4: [[ISNOTDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
-// CK4: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]]
+// CK4: [[CMP1:%.+]] = and i1 [[ISARRAY]], [[ISNOTDEL]]
// CK4: br i1 [[CMP1]], label %[[EVALDEL:[^,]+]], label %[[DONE]]
// CK4: [[EVALDEL]]
// CK4-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
// CK4-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
-// CK4-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
+// CK4-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], 281474976710652
// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]], {{.*}})
// CK4: br label %[[DONE]]
// CK4: [[DONE]]
diff --git a/clang/test/OpenMP/target_map_codegen_34.cpp b/clang/test/OpenMP/target_map_codegen_34.cpp
new file mode 100644
index 000000000000..43b07a033804
--- /dev/null
+++ b/clang/test/OpenMP/target_map_codegen_34.cpp
@@ -0,0 +1,258 @@
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+///==========================================================================///
+// RUN: %clang_cc1 -DCK34 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK34 --check-prefix CK34-64
+// RUN: %clang_cc1 -DCK34 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK34 --check-prefix CK34-64
+// RUN: %clang_cc1 -DCK34 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK34 --check-prefix CK34-32
+// RUN: %clang_cc1 -DCK34 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK34 --check-prefix CK34-32
+
+// RUN: %clang_cc1 -DCK34 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY32 %s
+// RUN: %clang_cc1 -DCK34 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY32 %s
+// RUN: %clang_cc1 -DCK34 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY32 %s
+// RUN: %clang_cc1 -DCK34 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY32 %s
+// SIMD-ONLY32-NOT: {{__kmpc|__tgt}}
+#ifdef CK34
+
+class C {
+public:
+ int a;
+ double *b;
+};
+
+#pragma omp declare mapper(C s) map(s.a, s.b[0:2])
+
+class S {
+ int a;
+ C c;
+ int b;
+public:
+ void foo();
+};
+
+// TARGET_PARAM = 0x20
+// MEMBER_OF_1 | TO = 0x1000000000001
+// MEMBER_OF_1 | IMPLICIT | TO = 0x1000000000201
+// CK34-DAG: [[MTYPE_TO:@.+]] = {{.+}}constant [4 x i64] [i64 [[#0x20]], i64 [[#0x1000000000001]], i64 [[#0x1000000000001]], i64 [[#0x1000000000201]]]
+// TARGET_PARAM = 0x20
+// MEMBER_OF_1 | FROM = 0x1000000000002
+// MEMBER_OF_1 | IMPLICIT | FROM = 0x1000000000202
+// CK34-DAG: [[MTYPE_FROM:@.+]] = {{.+}}constant [4 x i64] [i64 [[#0x20]], i64 [[#0x1000000000002]], i64 [[#0x1000000000002]], i64 [[#0x1000000000202]]]
+
+void default_mapper() {
+ S s;
+
+ // CK34-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 -1, i8* @{{.+}}, i32 4, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[4 x i{{.+}}]* [[MTYPE_TO]]{{.+}}, i8** null, i8** [[GEPMF:%.+]])
+ // CK34-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK34-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+ // CK34-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+ // CK34-DAG: [[GEPMF]] = bitcast [4 x i8*]* [[MF:%.+]] to i8**
+
+ // pass TARGET_PARAM {&s, &s, ((void*)(&s+1)-(void*)&s)}
+
+ // CK34-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK34-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK34-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+ // CK34-DAG: [[MF0:%.+]] = getelementptr inbounds {{.+}}[[MF]], i{{.+}} 0, i{{.+}} 0
+
+ // CK34-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to %class.S**
+ // CK34-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to %class.S**
+
+ // CK34-DAG: store %class.S* [[S_ADDR:%.+]], %class.S** [[BPC0]],
+ // CK34-DAG: store %class.S* [[S_ADDR]], %class.S** [[PC0]],
+ // CK34-DAG: store i64 [[S_SIZE:%.+]], i64* [[S0]],
+ // CK34-DAG: store i8* null, i8** [[MF0]],
+
+ // CK34-DAG: [[S_SIZE]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+ // CK34-DAG: [[SZ]] = sub i64 [[S_1_INTPTR:%.+]], [[S_INTPTR:%.+]]
+ // CK34-DAG: [[S_1_INTPTR]] = ptrtoint i8* [[S_1_VOID:%.+]] to i64
+ // CK34-DAG: [[S_INTPTR]] = ptrtoint i8* [[S_VOID:%.+]] to i64
+ // CK34-DAG: [[S_1_VOID]] = bitcast %class.S* [[S_1:%.+]] to i8*
+ // CK34-DAG: [[S_VOID]] = bitcast %class.S* [[S_ADDR]] to i8*
+ // CK34-DAG: [[S_1]] = getelementptr %class.S, %class.S* [[S_ADDR]], i32 1
+
+ // pass MEMBER_OF_1 | TO {&s, &s, ((void*)(&s.a+1)-(void*)&s)} to copy the data of s.a.
+
+ // CK34-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+ // CK34-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+ // CK34-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
+ // CK34-DAG: [[MF1:%.+]] = getelementptr inbounds {{.+}}[[MF]], i{{.+}} 0, i{{.+}} 1
+
+ // CK34-DAG: [[BPC1:%.+]] = bitcast i8** [[BP1]] to %class.S**
+ // CK34-DAG: [[PC1:%.+]] = bitcast i8** [[P1]] to %class.S**
+
+ // CK34-DAG: store %class.S* [[S_ADDR]], %class.S** [[BPC1]],
+ // CK34-DAG: store %class.S* [[S_ADDR]], %class.S** [[PC1]],
+ // CK34-DAG: store i64 [[A_SIZE:%.+]], i64* [[S1]],
+ // CK34-DAG: store i8* null, i8** [[MF1]],
+
+ // CK34-DAG: [[A_SIZE]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+ // CK34-DAG: [[SZ]] = sub i64 [[C_BEGIN_INTPTR:%.+]], [[S_INTPTR:%.+]]
+ // CK34-DAG: [[S_INTPTR]] = ptrtoint i8* [[S_VOID:%.+]] to i64
+ // CK34-DAG: [[C_BEGIN_INTPTR]] = ptrtoint i8* [[C_BEGIN_VOID:%.+]] to i64
+ // CK34-DAG: [[S_VOID]] = bitcast %class.S* [[S_ADDR]] to i8*
+ // CK34-DAG: [[C_BEGIN_VOID]] = bitcast %class.C* [[C_ADDR:%.+]] to i8*
+ // CK34-64-DAG: [[C_ADDR]] = getelementptr inbounds %class.S, %class.S* [[S_ADDR]], i32 0, i32 2
+ // CK34-32-DAG: [[C_ADDR]] = getelementptr inbounds %class.S, %class.S* [[S_ADDR]], i32 0, i32 1
+
+ // pass MEMBER_OF_1 | TO {&s, &s.c+1, ((void*)(&s)+31+1-(void*)(&s.c+1))} to copy the data of s.b.
+
+ // CK34-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
+ // CK34-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
+ // CK34-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2
+ // CK34-DAG: [[MF2:%.+]] = getelementptr inbounds {{.+}}[[MF]], i{{.+}} 0, i{{.+}} 2
+
+ // CK34-DAG: [[BPC2:%.+]] = bitcast i8** [[BP2]] to %class.S**
+ // CK34-DAG: [[PC2:%.+]] = bitcast i8** [[P2]] to %class.C**
+
+ // CK34-DAG: store %class.S* [[S_ADDR]], %class.S** [[BPC2]],
+ // CK34-DAG: store %class.C* [[C_END:%.+]], %class.C** [[PC2]],
+ // CK34-DAG: store i64 [[B_SIZE:%.+]], i64* [[S2]],
+ // CK34-DAG: store i8* null, i8** [[MF2]],
+
+ // CK34-DAG: [[C_END]] = getelementptr %class.C, %class.C* [[C_ADDR]], i{{.+}} 1
+
+ // CK34-DAG: [[B_SIZE]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+ // CK34-DAG: [[SZ]] = sub i64 [[S_END_INTPTR:%.+]], [[C_END_INTPTR:%.+]]
+ // CK34-DAG: [[C_END_INTPTR]] = ptrtoint i8* [[C_END_VOID:%.+]] to i64
+ // CK34-DAG: [[S_END_INTPTR]] = ptrtoint i8* [[S_END_VOID:%.+]] to i64
+ // CK34-DAG: [[C_END_VOID]] = bitcast %class.C* [[C_END]] to i8*
+ // CK34-DAG: [[S_END_VOID]] = getelementptr i8, i8* [[S_LAST:%.+]], i{{.+}} 1
+ // CK34-64-DAG: [[S_LAST]] = getelementptr i8, i8* [[S_VOID:%.+]], i64 31
+ // CK34-32-DAG: [[S_LAST]] = getelementptr i8, i8* [[S_VOID:%.+]], i32 15
+ // CK34-DAG: [[S_VOID]] = bitcast %class.S* [[S_ADDR]] to i8*
+
+ // pass MEMBER_OF_1 | TO | IMPLICIT | MAPPER {&s, &s.c, 16} to copy the data of s.c.
+
+ // CK34-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3
+ // CK34-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3
+ // CK34-DAG: [[S3:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 3
+ // CK34-DAG: [[MF3:%.+]] = getelementptr inbounds {{.+}}[[MF]], i{{.+}} 0, i{{.+}} 3
+
+ // CK34-DAG: [[BPC3:%.+]] = bitcast i8** [[BP3]] to %class.S**
+ // CK34-DAG: [[PC3:%.+]] = bitcast i8** [[P3]] to %class.C**
+
+ // CK34-DAG: store %class.S* [[S_ADDR]], %class.S** [[BPC3]],
+ // CK34-DAG: store %class.C* [[C_ADDR:%.+]], %class.C** [[PC3]],
+ // CK34-64-DAG: store i64 16, i64* [[S3]],
+ // CK34-32-DAG: store i64 8, i64* [[S3]],
+ // CK34-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64, i8*)* [[C_DEFAULT_MAPPER:@.+]] to i8*), i8** [[MF3]],
+
+ // CK34-64-DAG: [[C_ADDR]] = getelementptr inbounds %class.S, %class.S* [[S_ADDR]], i32 0, i32 2
+ // CK34-32-DAG: [[C_ADDR]] = getelementptr inbounds %class.S, %class.S* [[S_ADDR]], i32 0, i32 1
+
+ #pragma omp target map(to: s)
+ s.foo();
+
+ // CK34 : call void
+
+ // CK34-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 -1, i8* @{{.+}}, i32 4, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[4 x i{{.+}}]* [[MTYPE_FROM]]{{.+}}, i8** null, i8** [[GEPMF:%.+]])
+ // CK34-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK34-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+ // CK34-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+ // CK34-DAG: [[GEPMF]] = bitcast [4 x i8*]* [[MF:%.+]] to i8**
+
+ // pass TARGET_PARAM {&s, &s, ((void*)(&s+1)-(void*)&s)}
+
+ // CK34-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK34-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK34-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+ // CK34-DAG: [[MF0:%.+]] = getelementptr inbounds {{.+}}[[MF]], i{{.+}} 0, i{{.+}} 0
+
+ // CK34-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to %class.S**
+ // CK34-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to %class.S**
+
+ // CK34-DAG: store %class.S* [[S_ADDR]], %class.S** [[BPC0]],
+ // CK34-DAG: store %class.S* [[S_ADDR]], %class.S** [[PC0]],
+ // CK34-DAG: store i64 [[S_SIZE:%.+]], i64* [[S0]],
+ // CK34-DAG: store i8* null, i8** [[MF0]],
+
+ // CK34-DAG: [[S_SIZE]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+ // CK34-DAG: [[SZ]] = sub i64 [[S_1_INTPTR:%.+]], [[S_INTPTR:%.+]]
+ // CK34-DAG: [[S_1_INTPTR]] = ptrtoint i8* [[S_1_VOID:%.+]] to i64
+ // CK34-DAG: [[S_INTPTR]] = ptrtoint i8* [[S_VOID:%.+]] to i64
+ // CK34-DAG: [[S_1_VOID]] = bitcast %class.S* [[S_1:%.+]] to i8*
+ // CK34-DAG: [[S_VOID]] = bitcast %class.S* [[S_ADDR]] to i8*
+ // CK34-DAG: [[S_1]] = getelementptr %class.S, %class.S* [[S_ADDR]], i32 1
+
+ // pass MEMBER_OF_1 | FROM {&s, &s, ((void*)(&s.a+1)-(void*)&s)} to copy the data of s.a.
+
+ // CK34-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+ // CK34-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+ // CK34-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
+ // CK34-DAG: [[MF1:%.+]] = getelementptr inbounds {{.+}}[[MF]], i{{.+}} 0, i{{.+}} 1
+
+ // CK34-DAG: [[BPC1:%.+]] = bitcast i8** [[BP1]] to %class.S**
+ // CK34-DAG: [[PC1:%.+]] = bitcast i8** [[P1]] to %class.S**
+
+ // CK34-DAG: store %class.S* [[S_ADDR]], %class.S** [[BPC1]],
+ // CK34-DAG: store %class.S* [[S_ADDR]], %class.S** [[PC1]],
+ // CK34-DAG: store i64 [[A_SIZE:%.+]], i64* [[S1]],
+ // CK34-DAG: store i8* null, i8** [[MF1]],
+
+ // CK34-DAG: [[A_SIZE]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+ // CK34-DAG: [[SZ]] = sub i64 [[C_BEGIN_INTPTR:%.+]], [[S_INTPTR:%.+]]
+ // CK34-DAG: [[S_INTPTR]] = ptrtoint i8* [[S_VOID:%.+]] to i64
+ // CK34-DAG: [[C_BEGIN_INTPTR]] = ptrtoint i8* [[C_BEGIN_VOID:%.+]] to i64
+ // CK34-DAG: [[S_VOID]] = bitcast %class.S* [[S_ADDR]] to i8*
+ // CK34-DAG: [[C_BEGIN_VOID]] = bitcast %class.C* [[C_ADDR:%.+]] to i8*
+ // CK34-64-DAG: [[C_ADDR]] = getelementptr inbounds %class.S, %class.S* [[S_ADDR]], i32 0, i32 2
+ // CK34-32-DAG: [[C_ADDR]] = getelementptr inbounds %class.S, %class.S* [[S_ADDR]], i32 0, i32 1
+
+ // pass MEMBER_OF_1 | FROM {&s, &s.c+1, ((void*)(&s)+31+1-(void*)(&s.c+1))} to copy the data of s.b.
+
+ // CK34-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
+ // CK34-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
+ // CK34-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2
+ // CK34-DAG: [[MF2:%.+]] = getelementptr inbounds {{.+}}[[MF]], i{{.+}} 0, i{{.+}} 2
+
+ // CK34-DAG: [[BPC2:%.+]] = bitcast i8** [[BP2]] to %class.S**
+ // CK34-DAG: [[PC2:%.+]] = bitcast i8** [[P2]] to %class.C**
+
+ // CK34-DAG: store %class.S* [[S_ADDR]], %class.S** [[BPC2]],
+ // CK34-DAG: store %class.C* [[C_END:%.+]], %class.C** [[PC2]],
+ // CK34-DAG: store i64 [[B_SIZE:%.+]], i64* [[S2]],
+ // CK34-DAG: store i8* null, i8** [[MF2]],
+
+ // CK34-DAG: [[C_END]] = getelementptr %class.C, %class.C* [[C_ADDR]], i{{.+}} 1
+
+ // CK34-DAG: [[B_SIZE]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+ // CK34-DAG: [[SZ]] = sub i64 [[S_END_INTPTR:%.+]], [[C_END_INTPTR:%.+]]
+ // CK34-DAG: [[C_END_INTPTR]] = ptrtoint i8* [[C_END_VOID:%.+]] to i64
+ // CK34-DAG: [[S_END_INTPTR]] = ptrtoint i8* [[S_END_VOID:%.+]] to i64
+ // CK34-DAG: [[C_END_VOID]] = bitcast %class.C* [[C_END]] to i8*
+ // CK34-DAG: [[S_END_VOID]] = getelementptr i8, i8* [[S_LAST:%.+]], i{{.+}} 1
+ // CK34-64-DAG: [[S_LAST]] = getelementptr i8, i8* [[S_VOID:%.+]], i64 31
+ // CK34-32-DAG: [[S_LAST]] = getelementptr i8, i8* [[S_VOID:%.+]], i32 15
+ // CK34-DAG: [[S_VOID]] = bitcast %class.S* [[S_ADDR]] to i8*
+
+ // pass MEMBER_OF_1 | FROM | IMPLICIT | MAPPER {&s, &s.c, 16} to copy the data of s.c.
+
+ // CK34-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3
+ // CK34-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3
+ // CK34-DAG: [[S3:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 3
+ // CK34-DAG: [[MF3:%.+]] = getelementptr inbounds {{.+}}[[MF]], i{{.+}} 0, i{{.+}} 3
+
+ // CK34-DAG: [[BPC3:%.+]] = bitcast i8** [[BP3]] to %class.S**
+ // CK34-DAG: [[PC3:%.+]] = bitcast i8** [[P3]] to %class.C**
+
+ // CK34-DAG: store %class.S* [[S_ADDR]], %class.S** [[BPC3]],
+ // CK34-DAG: store %class.C* [[C_ADDR:%.+]], %class.C** [[PC3]],
+ // CK34-64-DAG: store i64 16, i64* [[S3]],
+ // CK34-32-DAG: store i64 8, i64* [[S3]],
+ // CK34-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64, i8*)* [[C_DEFAULT_MAPPER]] to i8*), i8** [[MF3]],
+
+ // CK34-64-DAG: [[C_ADDR]] = getelementptr inbounds %class.S, %class.S* [[S_ADDR]], i32 0, i32 2
+ // CK34-32-DAG: [[C_ADDR]] = getelementptr inbounds %class.S, %class.S* [[S_ADDR]], i32 0, i32 1
+
+ #pragma omp target map(from: s)
+ s.foo();
+}
+
+#endif // CK34
+#endif
diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp
index f273b6618e1a..6731439de9ce 100644
--- a/openmp/libomptarget/src/omptarget.cpp
+++ b/openmp/libomptarget/src/omptarget.cpp
@@ -269,10 +269,11 @@ int targetDataMapper(ident_t *loc, DeviceTy &Device, void *arg_base, void *arg,
MapperArgNames[I] = C.Name;
}
- int rc = target_data_function(
- loc, Device, MapperComponents.Components.size(), MapperArgsBase.data(),
- MapperArgs.data(), MapperArgSizes.data(), MapperArgTypes.data(),
- MapperArgNames.data(), /*arg_mappers*/ nullptr, AsyncInfo);
+ int rc = target_data_function(loc, Device, MapperComponents.Components.size(),
+ MapperArgsBase.data(), MapperArgs.data(),
+ MapperArgSizes.data(), MapperArgTypes.data(),
+ MapperArgNames.data(), /*arg_mappers*/ nullptr,
+ AsyncInfo, /*FromMapper=*/true);
return rc;
}
@@ -281,7 +282,8 @@ int targetDataMapper(ident_t *loc, DeviceTy &Device, void *arg_base, void *arg,
int targetDataBegin(ident_t *loc, DeviceTy &Device, int32_t arg_num,
void **args_base, void **args, int64_t *arg_sizes,
int64_t *arg_types, map_var_info_t *arg_names,
- void **arg_mappers, AsyncInfoTy &AsyncInfo) {
+ void **arg_mappers, AsyncInfoTy &AsyncInfo,
+ bool FromMapper) {
// process each input.
for (int32_t i = 0; i < arg_num; ++i) {
// Ignore private variables and arrays - there is no mapping for them.
@@ -379,7 +381,10 @@ int targetDataBegin(ident_t *loc, DeviceTy &Device, int32_t arg_num,
Pointer_HstPtrBegin = HstPtrBase;
// modify current entry.
HstPtrBase = *(void **)HstPtrBase;
- UpdateRef = true; // subsequently update ref count of pointee
+ // No need to update pointee ref count for the first element of the
+ // subelement that comes from mapper.
+ UpdateRef =
+ (!FromMapper || i != 0); // subsequently update ref count of pointee
}
void *TgtPtrBegin = Device.getOrAllocTgtPtr(
@@ -483,7 +488,7 @@ struct DeallocTgtPtrInfo {
int targetDataEnd(ident_t *loc, DeviceTy &Device, int32_t ArgNum,
void **ArgBases, void **Args, int64_t *ArgSizes,
int64_t *ArgTypes, map_var_info_t *ArgNames,
- void **ArgMappers, AsyncInfoTy &AsyncInfo) {
+ void **ArgMappers, AsyncInfoTy &AsyncInfo, bool FromMapper) {
int Ret;
std::vector<DeallocTgtPtrInfo> DeallocTgtPtrs;
// process each input.
@@ -536,7 +541,8 @@ int targetDataEnd(ident_t *loc, DeviceTy &Device, int32_t ArgNum,
bool IsLast, IsHostPtr;
bool IsImplicit = ArgTypes[I] & OMP_TGT_MAPTYPE_IMPLICIT;
bool UpdateRef = !(ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) ||
- (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ);
+ (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ &&
+ (!FromMapper || I != ArgNum - 1));
bool ForceDelete = ArgTypes[I] & OMP_TGT_MAPTYPE_DELETE;
bool HasCloseModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_CLOSE;
bool HasPresentModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_PRESENT;
@@ -584,8 +590,13 @@ int targetDataEnd(ident_t *loc, DeviceTy &Device, int32_t ArgNum,
bool DelEntry = IsLast || ForceDelete;
- if ((ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) &&
- !(ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)) {
+ // If the last element from the mapper (for end transfer args comes in
+ // reverse order), do not remove the partial entry, the parent struct still
+ // exists.
+ if (((ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) &&
+ !(ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)) ||
+ (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ && FromMapper &&
+ I == ArgNum - 1)) {
DelEntry = false; // protect parent struct from being deallocated
}
@@ -822,7 +833,7 @@ static int getNonContigMergedDimension(__tgt_target_non_contig *NonContig,
int targetDataUpdate(ident_t *loc, DeviceTy &Device, int32_t ArgNum,
void **ArgsBase, void **Args, int64_t *ArgSizes,
int64_t *ArgTypes, map_var_info_t *ArgNames,
- void **ArgMappers, AsyncInfoTy &AsyncInfo) {
+ void **ArgMappers, AsyncInfoTy &AsyncInfo, bool) {
// process each input.
for (int32_t I = 0; I < ArgNum; ++I) {
if ((ArgTypes[I] & OMP_TGT_MAPTYPE_LITERAL) ||
diff --git a/openmp/libomptarget/src/private.h b/openmp/libomptarget/src/private.h
index fc6997a2d977..2eb0c812e95c 100644
--- a/openmp/libomptarget/src/private.h
+++ b/openmp/libomptarget/src/private.h
@@ -23,17 +23,20 @@
extern int targetDataBegin(ident_t *loc, DeviceTy &Device, int32_t arg_num,
void **args_base, void **args, int64_t *arg_sizes,
int64_t *arg_types, map_var_info_t *arg_names,
- void **arg_mappers, AsyncInfoTy &AsyncInfo);
+ void **arg_mappers, AsyncInfoTy &AsyncInfo,
+ bool FromMapper = false);
extern int targetDataEnd(ident_t *loc, DeviceTy &Device, int32_t ArgNum,
void **ArgBases, void **Args, int64_t *ArgSizes,
int64_t *ArgTypes, map_var_info_t *arg_names,
- void **ArgMappers, AsyncInfoTy &AsyncInfo);
+ void **ArgMappers, AsyncInfoTy &AsyncInfo,
+ bool FromMapper = false);
extern int targetDataUpdate(ident_t *loc, DeviceTy &Device, int32_t arg_num,
void **args_base, void **args, int64_t *arg_sizes,
int64_t *arg_types, map_var_info_t *arg_names,
- void **arg_mappers, AsyncInfoTy &AsyncInfo);
+ void **arg_mappers, AsyncInfoTy &AsyncInfo,
+ bool FromMapper = false);
extern int target(ident_t *loc, DeviceTy &Device, void *HostPtr, int32_t ArgNum,
void **ArgBases, void **Args, int64_t *ArgSizes,
@@ -76,7 +79,8 @@ typedef void (*MapperFuncPtrTy)(void *, void *, void *, int64_t, int64_t,
// targetDataEnd and targetDataUpdate).
typedef int (*TargetDataFuncPtrTy)(ident_t *, DeviceTy &, int32_t, void **,
void **, int64_t *, int64_t *,
- map_var_info_t *, void **, AsyncInfoTy &);
+ map_var_info_t *, void **, AsyncInfoTy &,
+ bool);
// Implemented in libomp, they are called from within __tgt_* functions.
#ifdef __cplusplus
diff --git a/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers.cpp b/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers.cpp
new file mode 100644
index 000000000000..ae2902a3d08d
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers.cpp
@@ -0,0 +1,63 @@
+// RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
+
+#include <cstdio>
+#include <cstdlib>
+
+typedef struct {
+ int a;
+ double *b;
+} C1;
+#pragma omp declare mapper(C1 s) map(to : s.a) map(from : s.b [0:2])
+
+typedef struct {
+ int a;
+ double *b;
+ C1 c;
+} C;
+#pragma omp declare mapper(C s) map(to : s.a, s.c) map(from : s.b [0:2])
+
+typedef struct {
+ int e;
+ C f;
+ int h;
+} D;
+
+int main() {
+ constexpr int N = 10;
+ D s;
+ s.e = 111;
+ s.f.a = 222;
+ s.f.c.a = 777;
+ double x[2];
+ double x1[2];
+ x[1] = 20;
+ s.f.b = &x[0];
+ s.f.c.b = &x1[0];
+ s.h = N;
+
+ D *sp = &s;
+ D **spp = &sp;
+
+ printf("%d %d %d %4.5f %d\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.c.a,
+ spp[0][0].f.b[1], spp[0][0].f.b == &x[0] ? 1 : 0);
+ // CHECK: 111 222 777 20.00000 1
+
+ __intptr_t p = reinterpret_cast<__intptr_t>(&x[0]);
+#pragma omp target map(tofrom : spp[0][0]) firstprivate(p)
+ {
+ printf("%d %d %d\n", spp[0][0].f.a, spp[0][0].f.c.a,
+ spp[0][0].f.b == reinterpret_cast<void *>(p) ? 1 : 0);
+ // CHECK: 222 777 0
+ spp[0][0].e = 333;
+ spp[0][0].f.a = 444;
+ spp[0][0].f.c.a = 555;
+ spp[0][0].f.b[1] = 40;
+ }
+ printf("%d %d %d %4.5f %d\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.c.a,
+ spp[0][0].f.b[1], spp[0][0].f.b == &x[0] ? 1 : 0);
+ // CHECK: 333 222 777 40.00000 1
+}
More information about the cfe-commits
mailing list