[OPENMP50]Mapping of the subcomponents with the 'default' mappers.

[ABataev]

Summary

If the mapped structure has data members, which have 'default' mappers,
need to map these members individually using their 'default' mappers.

[abhinavgaba]

The patch currently asserts for the following:

#include <stdio.h>
                                                                
typedef struct { int a; double *b; } C;                         
#pragma omp declare mapper(default: C s) map(s.a, s.b[0:2])     

int main() {
C s;                                                            
s.a = 10;                                                       
double x[2]; x[0] = 20;                                         
s.b = &x[0];
C* sp = &s;                                                     
C** spp = &sp;

printf("%d %lf %p\n", spp[0][0].a, spp[0][0].b[0], spp[0][0].b);
                                                                
#pragma omp target map(tofrom:spp[0][0])
printf("%d %lf %p\n", spp[0][0].a, spp[0][0].b[0], spp[0][0].b);
}

Note that this doesn't work correctly before this patch either. This changes it from a runfail to a compfail.

[ABataev]

Temp update, not finished yet. Still WIP.

[ABataev]

Fixes and updates

[ABataev]

Removed unnecessary changes.

[grokos]

Libomptarget changes look good, I'll let @jdoerfert provide feedback for the clang part.

clang/include/clang/AST/OpenMPClause.h

const_component_lists_iterator &operator++() {
       } else {
         std::advance(this->I, *ListSizeCur - PrevListSize);
         PrevListSize = *ListSizeCur;
 
         // We are done with a declaration, move to the next one.
         if (!(--RemainingLists)) {
           ++DeclCur;
           ++NumListsCur;
-          if (SupportsMapper)
-            ++MapperCur;
           RemainingLists = *NumListsCur;
           assert(RemainingLists && "No lists in the following declaration??");
         }
       }
 
       ++ListSizeCur;
+      if (SupportsMapper)
+        ++MapperCur;
       return *this;
     }
   };
 
   using const_component_lists_range =
       llvm::iterator_range<const_component_lists_iterator>;
 
   /// Iterators for all component lists.

clang/lib/CodeGen/CGOpenMPRuntime.cpp

struct MapCombinedInfoTy {
     }
   };
 
   /// Map between a struct and the its lowest & highest elements which have been
   /// mapped.
   /// [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;
   };
 
 private:
   /// Kind that defines how a device pointer has to be returned.
   struct MapInfo {
     OMPClauseMappableExprCommon::MappableExprComponentListRef Components;
for (; I != CE; ++I) {
 
         // If this component is a pointer inside the base struct then we don't
         // need to create any entry for it - it will be combined with the object
         // it is pointing to into a single PTR_AND_OBJ entry.
         bool IsMemberPointerOrAddr =
             (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
           // device.
           PartialStruct.LowestElem = {0, LB};
           CharUnits TypeSize = CGF.getContext().getTypeSizeInChars(
               I->getAssociatedExpression()->getType());
           Address HB = CGF.Builder.CreateConstGEP(
               CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(LB,
                                                               CGF.VoidPtrTy),
               TypeSize.getQuantity() - 1);
           PartialStruct.HighestElem = {
               std::numeric_limits<decltype(
                   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
                    Component : OverlappedElements) {
             Address ComponentLB = Address::invalid();
             for (const OMPClauseMappableExprCommon::MappableComponent &MC :
                  Component) {
               if (MC.getAssociatedDeclaration()) {
for (; I != CE; ++I) {
               Address HB =
                   CGF.EmitOMPArraySectionExpr(OASE, /*IsLowerBound=*/false)
                       .getAddress(CGF);
               PartialStruct.HighestElem = {FieldIndex, HB};
             } else {
               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) {
             PartialStruct.HighestElem = {FieldIndex, LB};
           }
         }
 
         // Need to emit combined struct for array sections.
void emitCombinedEntry(MapCombinedInfoTy &CombinedInfo,
                          bool NotTargetParams = true) const {
     if (CurTypes.size() == 1 &&
         ((CurTypes.back() & OMP_MAP_MEMBER_OF) != OMP_MAP_MEMBER_OF) &&
         !PartialStruct.IsArraySection)
       return;
     Address LBAddr = PartialStruct.LowestElem.second;
     Address HBAddr = PartialStruct.HighestElem.second;
     if (PartialStruct.HasCompleteRecord) {
-      LBAddr = PartialStruct.Base;
-      HBAddr = PartialStruct.Base;
+      LBAddr = PartialStruct.LB;
+      HBAddr = PartialStruct.LB;
     }
     CombinedInfo.Exprs.push_back(VD);
     // Base is the base of the struct
     CombinedInfo.BasePointers.push_back(PartialStruct.Base.getPointer());
     // Pointer is the address of the lowest element
     llvm::Value *LB = LBAddr.getPointer();
     CombinedInfo.Pointers.push_back(LB);
     // There should not be a mapper for a combined entry.
for (const MapData &L : DeclComponentLists) {
           auto &OverlappedElements = OverlappedData.FindAndConstruct(&BaseData);
           OverlappedElements.getSecond().push_back(SubData);
         }
       }
     }
     // 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());
       }
     }
     for (auto &Pair : OverlappedData) {
       llvm::stable_sort(
           Pair.getSecond(),
           [&Layout](
               OMPClauseMappableExprCommon::MappableExprComponentListRef First,

 
 /// Emit the user-defined mapper function. The code generation follows the
 /// pattern in the example below.
 /// \code
 /// void .omp_mapper.<type_name>.<mapper_id>.(void *rt_mapper_handle,
 ///                                           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:
 ///     for (auto c : begin[i]->all_components) {
 ///       if (c.hasMapper())
 ///         (*c.Mapper())(rt_mapper_handle, c.arg_base, c.arg_begin, c.arg_size,
 ///                       c.arg_type, c.arg_name);
 ///       else
 ///         __tgt_push_mapper_component(rt_mapper_handle, c.arg_base,
 ///                                     c.arg_begin, c.arg_size, c.arg_type,
 ///                                     c.arg_name);
 ///     }
 ///   }
 ///   // 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,
                                             CodeGenFunction *CGF) {
   if (UDMMap.count(D) > 0)
     return;
   ASTContext &C = CGM.getContext();
   QualType Ty = D->getType();
void CGOpenMPRuntime::emitUDMapperArrayInitOrDel(
     CharUnits ElementSize, llvm::BasicBlock *ExitBB, bool IsInit) {
   StringRef Prefix = IsInit ? ".init" : ".del";
 
   // Evaluate if this is an array section.
   llvm::BasicBlock *BodyBB =
       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"}));
   }
   Cond = MapperCGF.Builder.CreateAnd(Cond, DeleteCond);
   MapperCGF.Builder.CreateCondBr(Cond, BodyBB, ExitBB);
 
   MapperCGF.EmitBlock(BodyBB);
   // Get the array size by multiplying element size and element number (i.e., \p
   // Size).
   llvm::Value *ArraySize = MapperCGF.Builder.CreateNUWMul(
       Size, MapperCGF.Builder.getInt64(ElementSize.getQuantity()));
   // Remove OMP_MAP_TO and OMP_MAP_FROM from the map type, so that it achieves
   // memory allocation/deletion purpose only.
   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
   // data structure.
   llvm::Value *OffloadingArgs[] = {Handle,    Base,       Begin,
                                    ArraySize, MapTypeArg, MapNameArg};
   MapperCGF.EmitRuntimeCall(
       OMPBuilder.getOrCreateRuntimeFunction(CGM.getModule(),
for (CapturedStmt::const_capture_iterator CI = CS.capture_begin(),
       assert(CurInfo.BasePointers.size() == CurInfo.Pointers.size() &&
              CurInfo.BasePointers.size() == CurInfo.Sizes.size() &&
              CurInfo.BasePointers.size() == CurInfo.Types.size() &&
              CurInfo.BasePointers.size() == CurInfo.Mappers.size() &&
              "Inconsistent map information sizes!");
 
       // 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);
     }
     // Adjust MEMBER_OF flags for the lambdas captures.
     MEHandler.adjustMemberOfForLambdaCaptures(
         LambdaPointers, CombinedInfo.BasePointers, CombinedInfo.Pointers,
         CombinedInfo.Types);

clang/lib/Sema/SemaOpenMP.cpp

for (Expr *E : AC->varlists()) {
                                             AllocatorKind, AC->getAllocator()))
         continue;
       applyOMPAllocateAttribute(S, PrivateVD, AllocatorKind, AC->getAllocator(),
                                 E->getSourceRange());
     }
   }
 }
 
+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,
     Stmt *AStmt, SourceLocation StartLoc, SourceLocation EndLoc) {
   StmtResult Res = StmtError();
   // First check CancelRegion which is then used in checkNestingOfRegions.
   if (checkCancelRegion(*this, Kind, CancelRegion, StartLoc) ||
       checkNestingOfRegions(*this, DSAStack, Kind, DirName, CancelRegion,
for (unsigned I = 0, E = DefaultmapKindNum; I < E; ++I) {
           ClausesWithImplicit.emplace_back(Implicit);
           ErrorFound |= cast<OMPMapClause>(Implicit)->varlist_size() !=
                         ImplicitMap.size();
         } else {
           ErrorFound = true;
         }
       }
     }
+    // Build expressions for implicit maps of data members with 'default'
+    // mappers.
+    if (LangOpts.OpenMP >= 50)
+      processImplicitMapsWithDefaultMappers(*this, DSAStack,
+                                            ClausesWithImplicit);
   }
 
   llvm::SmallVector<OpenMPDirectiveKind, 4> AllowedNameModifiers;
   switch (Kind) {
   case OMPD_parallel:
     Res = ActOnOpenMPParallelDirective(ClausesWithImplicit, AStmt, StartLoc,
                                        EndLoc);
     AllowedNameModifiers.push_back(OMPD_parallel);
bool VisitCXXThisExpr(CXXThisExpr *CTE) {
     Components.emplace_back(CTE, nullptr, IsNonContiguous);
     return true;
   }
   bool VisitCXXOperatorCallExpr(CXXOperatorCallExpr *COCE) {
     assert(!RelevantExpr && "RelevantExpr is expected to be nullptr");
     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;
   }
   const Expr *getFoundBase() const {
     return RelevantExpr;
   }
   explicit MapBaseChecker(
Sema::DeclGroupPtrTy Sema::ActOnOpenMPDeclareMapperDirective(
   const auto I = PreviousRedeclTypes.find(MapperType.getCanonicalType());
   bool Invalid = false;
   if (I != PreviousRedeclTypes.end()) {
     Diag(StartLoc, diag::err_omp_declare_mapper_redefinition)
         << MapperType << Name;
     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
     DC->addDecl(DMD);
   DMD->setAccess(AS);
   if (Invalid)
     DMD->setInvalidDecl();
 

clang/test/OpenMP/declare_mapper_codegen.cpp

 // CK0-DAG: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1
 // CK0-DAG: [[PTRBEGIN:%.+]] = bitcast i8* [[BEGIN]] to %class.C*
 // CK0-DAG: [[PTREND:%.+]] = getelementptr %class.C, %class.C* [[PTRBEGIN]], i64 [[SIZE]]
 // CK0-DAG: [[BPTRI:%.+]] = ptrtoint i8* [[BPTR]] to i64
 // CK0-DAG: [[PTRI:%.+]] = ptrtoint i8* [[BEGIN]] to i64
 // CK0-DAG: [[DIF:%.+]] = sub i64 [[BPTRI]], [[PTRI]]
 // CK0-DAG: [[NORM:%.+]] = sdiv exact i64 [[DIF]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
 // CK0-DAG: [[PTRSNE:%.+]] = icmp ne i64 [[NORM]], 0
-// CK0-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[PTRSNE]]
+// CK0-DAG: [[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]]
 // CK0: br i1 [[CMP1]], label %[[INIT:[^,]+]], label %[[LHEAD:[^,]+]]
 // CK0: [[INIT]]
 // CK0-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
 // CK0-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
-// 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:[^,]+]]
 
 // 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]], %{{.+}} ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]

 // CK0-DAG: [[TYPE2:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
 // CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR2BC]], i8* [[PTRADDR2BC]], i64 16, i64 [[TYPE2]], {{.*}})
 // CK0: [[PTRNEXT]] = getelementptr %class.C, %class.C* [[PTR]], i32 1
 // CK0: [[ISDONE:%.+]] = icmp eq %class.C* [[PTRNEXT]], [[PTREND]]
 // CK0: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
 
 // 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]]
 // CK0: ret void
 
 
 // CK0-LABEL: define {{.*}}void @{{.*}}foo{{.*}}
 void foo(int a){

 // CK1-DAG: [[PTRBEGIN:%.+]] = bitcast i8* [[BEGIN]] to %class.C*
 // CK1-DAG: [[PTREND:%.+]] = getelementptr %class.C, %class.C* [[PTRBEGIN]], i64 [[SIZE]]
 // CK1-DAG: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1
 // CK1-DAG: [[BPTRI:%.+]] = ptrtoint i8* [[BPTR]] to i64
 // CK1-DAG: [[PTRI:%.+]] = ptrtoint i8* [[BEGIN]] to i64
 // CK1-DAG: [[DIF:%.+]] = sub i64 [[BPTRI]], [[PTRI]]
 // CK1-DAG: [[NORM:%.+]] = sdiv exact i64 [[DIF]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
 // CK1-DAG: [[PTRSNE:%.+]] = icmp ne i64 [[NORM]], 0
-// CK1-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[PTRSNE]]
+// CK1-DAG: [[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]]
 // CK1: br i1 [[CMP1]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]]
 
 // 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:[^,]+]]
 
 // 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]], %{{.+}} ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]

 // CK1-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
 // CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR1BC]], i8* [[PTRADDR1BC]], i64 4, i64 [[TYPE1]], {{.*}})
 // CK1: [[PTRNEXT]] = getelementptr %class.C, %class.C* [[PTR]], i32 1
 // CK1: [[ISDONE:%.+]] = icmp eq %class.C* [[PTRNEXT]], [[PTREND]]
 // CK1: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
 
 // 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]]
 // CK1: ret void
 
 #endif // CK1
 
 

 // CK2-DAG: [[PTRBEGIN:%.+]] = bitcast i8* [[BEGIN]] to %class.C*
 // CK2-DAG: [[PTREND:%.+]] = getelementptr %class.C, %class.C* [[PTRBEGIN]], i64 [[SIZE]]
 // CK2-DAG: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1
 // CK2-DAG: [[BPTRI:%.+]] = ptrtoint i8* [[BPTR]] to i64
 // CK2-DAG: [[PTRI:%.+]] = ptrtoint i8* [[BEGIN]] to i64
 // CK2-DAG: [[DIF:%.+]] = sub i64 [[BPTRI]], [[PTRI]]
 // CK2-DAG: [[NORM:%.+]] = sdiv exact i64 [[DIF]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
 // CK2-DAG: [[PTRSNE:%.+]] = icmp ne i64 [[NORM]], 0
-// CK2-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[PTRSNE]]
+// CK2-DAG: [[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]]
 // CK2: br i1 [[CMP1]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]]
 
 // 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:[^,]+]]
 
 // 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]], %{{.+}} ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]

 // CK2-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
 // CK2: call void [[BMPRFUNC]](i8* [[HANDLE]], i8* [[BPTRADDR1BC]], i8* [[PTRADDR1BC]], i64 8, i64 [[TYPE1]], {{.*}})
 // CK2: [[PTRNEXT]] = getelementptr %class.C, %class.C* [[PTR]], i32 1
 // CK2: [[ISDONE:%.+]] = icmp eq %class.C* [[PTRNEXT]], [[PTREND]]
 // CK2: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
 
 // 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]]
 // CK2: ret void
 
 #endif // CK2
 
 

 // CK4-DAG: [[PTRBEGIN:%.+]] = bitcast i8* [[BEGIN]] to %class.C*
 // CK4-DAG: [[PTREND:%.+]] = getelementptr %class.C, %class.C* [[PTRBEGIN]], i64 [[SIZE]]
 // CK4-DAG: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1
 // CK4-DAG: [[BPTRI:%.+]] = ptrtoint i8* [[BPTR]] to i64
 // CK4-DAG: [[PTRI:%.+]] = ptrtoint i8* [[BEGIN]] to i64
 // CK4-DAG: [[DIF:%.+]] = sub i64 [[BPTRI]], [[PTRI]]
 // CK4-DAG: [[NORM:%.+]] = sdiv exact i64 [[DIF]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
 // CK4-DAG: [[PTRSNE:%.+]] = icmp ne i64 [[NORM]], 0
-// CK4-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[PTRSNE]]
+// CK4-DAG: [[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]]
 // CK4: br i1 [[CMP1]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]]
 
 // 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:[^,]+]]
 
 // 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]], %{{.+}} ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]

 // CK4-DAG: [[TYPE2:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
 // CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR2BC]], i8* [[PTRADDR2BC]], i64 16, i64 [[TYPE2]], {{.*}})
 // CK4: [[PTRNEXT]] = getelementptr %class.C, %class.C* [[PTR]], i32 1
 // CK4: [[ISDONE:%.+]] = icmp eq %class.C* [[PTRNEXT]], [[PTREND]]
 // CK4: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
 
 // 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]]
 // CK4: ret void
 
 
 // CK4-LABEL: define {{.*}}void @{{.*}}foo{{.*}}
 void foo(int a){

clang/test/OpenMP/target_map_codegen_34.cpp

+// 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

openmp/libomptarget/src/omptarget.cpp

auto &C =
             .Components[target_data_function == targetDataEnd ? E - I - 1 : I];
     MapperArgsBase[I] = C.Base;
     MapperArgs[I] = C.Begin;
     MapperArgSizes[I] = C.Size;
     MapperArgTypes[I] = C.Type;
     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;
 }
 
 /// Internal function to do the mapping and transfer the data to the device
 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.
     if ((arg_types[i] & OMP_TGT_MAPTYPE_LITERAL) ||
         (arg_types[i] & OMP_TGT_MAPTYPE_PRIVATE))
       continue;
 
     if (arg_mappers && arg_mappers[i]) {
if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ) {
       }
       DP("There are %zu bytes allocated at target address " DPxMOD " - is%s new"
          "\n",
          sizeof(void *), DPxPTR(PointerTgtPtrBegin),
          (Pointer_IsNew ? "" : " not"));
       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(
         HstPtrBegin, HstPtrBase, data_size, HstPtrName, IsNew, IsHostPtr,
         IsImplicit, UpdateRef, HasCloseModifier, HasPresentModifier);
     // If data_size==0, then the argument could be a zero-length pointer to
     // NULL, so getOrAlloc() returning NULL is not an error.
     if (!TgtPtrBegin && (data_size || HasPresentModifier)) {
DeallocTgtPtrInfo(void *HstPtr, int64_t Size, bool ForceDelete,
         HasCloseModifier(HasCloseModifier) {}
 };
 } // namespace
 
 /// Internal function to undo the mapping and retrieve the data from the device.
 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.
   for (int32_t I = ArgNum - 1; I >= 0; --I) {
     // Ignore private variables and arrays - there is no mapping for them.
     // Also, ignore the use_device_ptr directive, it has no effect here.
     if ((ArgTypes[I] & OMP_TGT_MAPTYPE_LITERAL) ||
         (ArgTypes[I] & OMP_TGT_MAPTYPE_PRIVATE))
if (getParentIndex(ArgTypes[I]) < 0 && NextI < ArgNum &&
         HstPtrBegin = (char *)HstPtrBegin - Padding;
         DataSize += Padding;
       }
     }
 
     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;
 
     // If PTR_AND_OBJ, HstPtrBegin is address of pointee
     void *TgtPtrBegin = Device.getTgtPtrBegin(
         HstPtrBegin, DataSize, IsLast, UpdateRef, IsHostPtr, !IsImplicit);
     if (!TgtPtrBegin && (DataSize || HasPresentModifier)) {
for (int32_t I = ArgNum - 1; I >= 0; --I) {
     // construct and a corresponding list item of the original list item is not
     // present in the device data environment on exit from the region then the
     // list item is ignored."
     if (!TgtPtrBegin)
       continue;
 
     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
     }
 
     if ((ArgTypes[I] & OMP_TGT_MAPTYPE_FROM) || DelEntry) {
       // Move data back to the host
       if (ArgTypes[I] & OMP_TGT_MAPTYPE_FROM) {
         bool Always = ArgTypes[I] & OMP_TGT_MAPTYPE_ALWAYS;
         bool CopyMember = false;
static int getNonContigMergedDimension(__tgt_target_non_contig *NonContig,
   }
   return RemovedDim;
 }
 
 /// Internal function to pass data to/from the target.
 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) ||
         (ArgTypes[I] & OMP_TGT_MAPTYPE_PRIVATE))
       continue;
 
     if (ArgMappers && ArgMappers[I]) {
       // Instead of executing the regular path of targetDataUpdate, call the

openmp/libomptarget/src/private.h

 #include <SourceInfo.h>
 #include <omptarget.h>
 
 #include <cstdint>
 
 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,
                   int64_t *ArgTypes, map_var_info_t *arg_names,
                   void **ArgMappers, int32_t TeamNum, int32_t ThreadLimit,
                   int IsTeamConstruct, AsyncInfoTy &AsyncInfo);
 
 extern int CheckDeviceAndCtors(int64_t device_id);

 //                                           void * name);
 typedef void (*MapperFuncPtrTy)(void *, void *, void *, int64_t, int64_t,
                                 void *);
 
 // Function pointer type for targetData* functions (targetDataBegin,
 // 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
 extern "C" {
 #endif
 // functions that extract info from libomp; keep in sync
 int omp_get_default_device(void) __attribute__((weak));
 int32_t __kmpc_omp_taskwait(void *loc_ref, int32_t gtid) __attribute__((weak));

openmp/libomptarget/test/mapping/declare_mapper_nested_default_mappers.cpp

+// 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
+}