[OPENMP]Map data field with l-value reference types.

[ABataev]

Summary

Added initial support dfor the mapping of the data members with l-value
reference types.

[ABataev]

Rebase + fixes for references to complex data structures mapping

[ABataev]

Rebase + fixes

[jyu2]

Thanks Alexey for adding this! The changes LGTM, but I'm not super well-versed with this. You may wait a bit for other reviewers to weigh in before landing.

clang/lib/CodeGen/CGOpenMPRuntime.cpp

void generateInfoForComponentList(
     //   float f[50];
     // }
     // struct S2 {
     //   int i;
     //   float f[50];
     //   S1 s;
     //   double *p;
     //   struct S2 *ps;
+    //   int &ref;
     // }
     // S2 s;
     // S2 *ps;
     //
     // map(d)
     // &d, &d, sizeof(double), TARGET_PARAM | TO | FROM
     //
     // map(i)
void generateInfoForComponentList(
     // &s, &(s.p), sizeof(double*), MEMBER_OF(1) (**)
     // &(s.p), &(s.p[0]), 22*sizeof(double),
     //   MEMBER_OF(1) | PTR_AND_OBJ | TO (***)
     // (*) alloc space for struct members, only this is a target parameter
     // (**) map the pointer (nothing to be mapped in this example) (the compiler
     //      optimizes this entry out, same in the examples below)
     // (***) map the pointee (map: to)
     //
+    // map(to: s.ref)
+    // &s, &(s.ref), sizeof(int*), TARGET_PARAM (*)
+    // &s, &(s.ref), sizeof(int), MEMBER_OF(1) | PTR_AND_OBJ | TO (***)
+    // (*) alloc space for struct members, only this is a target parameter
+    // (**) map the pointer (nothing to be mapped in this example) (the compiler
+    //      optimizes this entry out, same in the examples below)
+    // (***) map the pointee (map: to)
+    //
     // map(s.ps)
     // &s, &(s.ps), sizeof(S2*), TARGET_PARAM | TO | FROM
     //
     // map(from: s.ps->s.i)
     // &s, &(s.ps), sizeof(S2*), TARGET_PARAM
     // &s, &(s.ps), sizeof(S2*), MEMBER_OF(1)
     // &(s.ps), &(s.ps->s.i), sizeof(int), MEMBER_OF(1) | PTR_AND_OBJ  | FROM
     //
void generateInfoForComponentList(
     // whether we are dealing with a member of a declared struct.
     const MemberExpr *EncounteredME = nullptr;
 
     // Track for the total number of dimension. Start from one for the dummy
     // dimension.
     uint64_t DimSize = 1;
 
     bool IsNonContiguous = CombinedInfo.NonContigInfo.IsNonContiguous;
+    bool IsPrevMemberReference = false;
 
     for (; I != CE; ++I) {
       // If the current component is member of a struct (parent struct) mark it.
       if (!EncounteredME) {
         EncounteredME = dyn_cast<MemberExpr>(I->getAssociatedExpression());
         // If we encounter a PTR_AND_OBJ entry from now on it should be marked
         // as MEMBER_OF the parent struct.
         if (EncounteredME) {
for (; I != CE; ++I) {
       const auto *UO = dyn_cast<UnaryOperator>(I->getAssociatedExpression());
       const auto *BO = dyn_cast<BinaryOperator>(I->getAssociatedExpression());
       bool IsPointer =
           OAShE ||
           (OASE && OMPArraySectionExpr::getBaseOriginalType(OASE)
                        .getCanonicalType()
                        ->isAnyPointerType()) ||
           I->getAssociatedExpression()->getType()->isAnyPointerType();
+      bool IsMemberReference = isa<MemberExpr>(I->getAssociatedExpression()) &&
+                               MapDecl &&
+                               MapDecl->getType()->isLValueReferenceType();
       bool IsNonDerefPointer = IsPointer && !UO && !BO && !IsNonContiguous;
 
       if (OASE)
         ++DimSize;
 
-      if (Next == CE || IsNonDerefPointer || IsFinalArraySection) {
+      if (Next == CE || IsMemberReference || IsNonDerefPointer ||
+          IsFinalArraySection) {
         // If this is not the last component, we expect the pointer to be
         // associated with an array expression or member expression.
         assert((Next == CE ||
                 isa<MemberExpr>(Next->getAssociatedExpression()) ||
                 isa<ArraySubscriptExpr>(Next->getAssociatedExpression()) ||
                 isa<OMPArraySectionExpr>(Next->getAssociatedExpression()) ||
                 isa<OMPArrayShapingExpr>(Next->getAssociatedExpression()) ||
                 isa<UnaryOperator>(Next->getAssociatedExpression()) ||
                 isa<BinaryOperator>(Next->getAssociatedExpression())) &&
                "Unexpected expression");
 
         Address LB = Address::invalid();
+        Address LowestElem = Address::invalid();
+        auto &&EmitMemberExprBase = [](CodeGenFunction &CGF,
+                                       const MemberExpr *E) {
+          const Expr *BaseExpr = E->getBase();
+          // If this is s.x, emit s as an lvalue.  If it is s->x, emit s as a
+          // scalar.
+          LValue BaseLV;
+          if (E->isArrow()) {
+            LValueBaseInfo BaseInfo;
+            TBAAAccessInfo TBAAInfo;
+            Address Addr =
+                CGF.EmitPointerWithAlignment(BaseExpr, &BaseInfo, &TBAAInfo);
+            QualType PtrTy = BaseExpr->getType()->getPointeeType();
+            BaseLV = CGF.MakeAddrLValue(Addr, PtrTy, BaseInfo, TBAAInfo);
+          } else {
+            BaseLV = CGF.EmitOMPSharedLValue(BaseExpr);
+          }
+          return BaseLV;
+        };
         if (OAShE) {
-          LB = Address(CGF.EmitScalarExpr(OAShE->getBase()),
+          LowestElem = LB = Address(CGF.EmitScalarExpr(OAShE->getBase()),
                                     CGF.getContext().getTypeAlignInChars(
                                         OAShE->getBase()->getType()));
+        } else if (IsMemberReference) {
+          const auto *ME = cast<MemberExpr>(I->getAssociatedExpression());
+          LValue BaseLVal = EmitMemberExprBase(CGF, ME);
+          LowestElem = CGF.EmitLValueForFieldInitialization(
+                              BaseLVal, cast<FieldDecl>(MapDecl))
+                           .getAddress(CGF);
+          LB = CGF.EmitLoadOfReferenceLValue(LowestElem, MapDecl->getType())
+                   .getAddress(CGF);
         } else {
-          LB = CGF.EmitOMPSharedLValue(I->getAssociatedExpression())
-                   .getAddress(CGF);
+          LowestElem = LB =
+              CGF.EmitOMPSharedLValue(I->getAssociatedExpression())
         }
 
         // 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);
+            EncounteredME &&
+            (((IsPointer || ForDeviceAddr) &&
+              I->getAssociatedExpression() == EncounteredME) ||
+             (IsPrevMemberReference && !IsPointer) ||
+             (IsMemberReference && Next != CE &&
+              !Next->getAssociatedExpression()->getType()->isPointerType()));
         if (!OverlappedElements.empty() && Next == CE) {
           // Handle base element with the info for overlapped elements.
           assert(!PartialStruct.Base.isValid() && "The base element is set.");
           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};
+          PartialStruct.LowestElem = {0, LowestElem};
           CharUnits TypeSize = CGF.getContext().getTypeSizeInChars(
               I->getAssociatedExpression()->getType());
           Address HB = CGF.Builder.CreateConstGEP(
-              CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(LB,
+              CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(LowestElem,
                                                               CGF.VoidPtrTy),
               TypeSize.getQuantity() - 1);
           PartialStruct.HighestElem = {
               std::numeric_limits<decltype(
                   PartialStruct.HighestElem.first)>::max(),
               HB};
           PartialStruct.Base = BP;
           PartialStruct.LB = LB;
for (; I != CE; ++I) {
                              /*AddIsTargetParamFlag=*/false, IsNonContiguous);
           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()) {
+              if (const ValueDecl *VD = MC.getAssociatedDeclaration()) {
+                const auto *FD = dyn_cast<FieldDecl>(VD);
+                if (FD && FD->getType()->isLValueReferenceType()) {
+                  const auto *ME =
+                      cast<MemberExpr>(MC.getAssociatedExpression());
+                  LValue BaseLVal = EmitMemberExprBase(CGF, ME);
+                  ComponentLB =
+                      CGF.EmitLValueForFieldInitialization(BaseLVal, FD)
+                          .getAddress(CGF);
+                } else {
                   ComponentLB =
                       CGF.EmitOMPSharedLValue(MC.getAssociatedExpression())
                           .getAddress(CGF);
+                }
                 Size = CGF.Builder.CreatePtrDiff(
                     CGF.EmitCastToVoidPtr(ComponentLB.getPointer()),
                     CGF.EmitCastToVoidPtr(LB.getPointer()));
                 break;
               }
             }
             assert(Size && "Failed to determine structure size");
             CombinedInfo.Exprs.emplace_back(MapDecl, MapExpr);
             CombinedInfo.BasePointers.push_back(BP.getPointer());
             CombinedInfo.Pointers.push_back(LB.getPointer());
             CombinedInfo.Sizes.push_back(CGF.Builder.CreateIntCast(
                 Size, CGF.Int64Ty, /*isSigned=*/true));
             CombinedInfo.Types.push_back(Flags);
             CombinedInfo.Mappers.push_back(nullptr);
             CombinedInfo.NonContigInfo.Dims.push_back(IsNonContiguous ? DimSize
                                                                       : 1);
             LB = CGF.Builder.CreateConstGEP(ComponentLB, 1);
           }
           CombinedInfo.Exprs.emplace_back(MapDecl, MapExpr);
           CombinedInfo.BasePointers.push_back(BP.getPointer());
           CombinedInfo.Pointers.push_back(LB.getPointer());
           Size = CGF.Builder.CreatePtrDiff(
-              CGF.EmitCastToVoidPtr(
-                  CGF.Builder.CreateConstGEP(HB, 1).getPointer()),
+              CGF.Builder.CreateConstGEP(HB, 1).getPointer(),
               CGF.EmitCastToVoidPtr(LB.getPointer()));
           CombinedInfo.Sizes.push_back(
               CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/true));
           CombinedInfo.Types.push_back(Flags);
           CombinedInfo.Mappers.push_back(nullptr);
           CombinedInfo.NonContigInfo.Dims.push_back(IsNonContiguous ? DimSize
                                                                     : 1);
           break;
for (; I != CE; ++I) {
 
           // We need to add a pointer flag for each map that comes from the
           // same expression except for the first one. We also need to signal
           // this map is the first one that relates with the current capture
           // (there is a set of entries for each capture).
           OpenMPOffloadMappingFlags Flags = getMapTypeBits(
               MapType, MapModifiers, MotionModifiers, IsImplicit,
               !IsExpressionFirstInfo || RequiresReference ||
-                  FirstPointerInComplexData,
+                  FirstPointerInComplexData || IsMemberReference,
               IsCaptureFirstInfo && !RequiresReference, IsNonContiguous);
 
-          if (!IsExpressionFirstInfo) {
+          if (!IsExpressionFirstInfo || IsMemberReference) {
             // If we have a PTR_AND_OBJ pair where the OBJ is a pointer as well,
             // then we reset the TO/FROM/ALWAYS/DELETE/CLOSE flags.
-            if (IsPointer)
+            if (IsPointer || (IsMemberReference && Next != CE))
               Flags &= ~(OMP_MAP_TO | OMP_MAP_FROM | OMP_MAP_ALWAYS |
                          OMP_MAP_DELETE | OMP_MAP_CLOSE);
 
             if (ShouldBeMemberOf) {
               // Set placeholder value MEMBER_OF=FFFF to indicate that the flag
               // should be later updated with the correct value of MEMBER_OF.
               Flags |= OMP_MAP_MEMBER_OF;
               // From now on, all subsequent PTR_AND_OBJ entries should not be
for (; I != CE; ++I) {
         // mapped member. If the parent is "*this", then the value declaration
         // is nullptr.
         if (EncounteredME) {
           const auto *FD = cast<FieldDecl>(EncounteredME->getMemberDecl());
           unsigned FieldIndex = FD->getFieldIndex();
 
           // Update info about the lowest and highest elements for this struct
           if (!PartialStruct.Base.isValid()) {
-            PartialStruct.LowestElem = {FieldIndex, LB};
+            PartialStruct.LowestElem = {FieldIndex, LowestElem};
             if (IsFinalArraySection) {
               Address HB =
                   CGF.EmitOMPArraySectionExpr(OASE, /*IsLowerBound=*/false)
                       .getAddress(CGF);
               PartialStruct.HighestElem = {FieldIndex, HB};
             } else {
-              PartialStruct.HighestElem = {FieldIndex, LB};
+              PartialStruct.HighestElem = {FieldIndex, LowestElem};
             }
             PartialStruct.Base = BP;
             PartialStruct.LB = BP;
           } else if (FieldIndex < PartialStruct.LowestElem.first) {
-            PartialStruct.LowestElem = {FieldIndex, LB};
+            PartialStruct.LowestElem = {FieldIndex, LowestElem};
           } else if (FieldIndex > PartialStruct.HighestElem.first) {
-            PartialStruct.HighestElem = {FieldIndex, LB};
+            PartialStruct.HighestElem = {FieldIndex, LowestElem};
           }
         }
 
         // Need to emit combined struct for array sections.
         if (IsFinalArraySection || IsNonContiguous)
           PartialStruct.IsArraySection = true;
 
         // If we have a final array section, we are done with this expression.
         if (IsFinalArraySection)
           break;
 
         // The pointer becomes the base for the next element.
         if (Next != CE)
-          BP = LB;
+          BP = IsMemberReference ? LowestElem : LB;
 
         IsExpressionFirstInfo = false;
         IsCaptureFirstInfo = false;
         FirstPointerInComplexData = false;
+        IsPrevMemberReference = IsMemberReference;
       } else if (FirstPointerInComplexData) {
         QualType Ty = Components.rbegin()
                           ->getAssociatedDeclaration()
                           ->getType()
                           .getNonReferenceType();
         BP = CGF.EmitLoadOfPointer(BP, Ty->castAs<PointerType>());
         FirstPointerInComplexData = false;
       }
for (const MapData &L : DeclComponentLists) {
           const auto It = (SI == SE) ? CI : SI;
           // If one component is a pointer and another one is a kind of
           // dereference of this pointer (array subscript, section, dereference,
           // etc.), it is not an overlapping.
           if (!isa<MemberExpr>(It->getAssociatedExpression()) ||
               std::prev(It)
                   ->getAssociatedExpression()
                   ->getType()
-                  .getNonReferenceType()
                   ->isPointerType())
             continue;
           const MapData &BaseData = CI == CE ? L : L1;
           OMPClauseMappableExprCommon::MappableExprComponentListRef SubData =
               SI == SE ? Components : Components1;
           auto &OverlappedElements = OverlappedData.FindAndConstruct(&BaseData);
           OverlappedElements.getSecond().push_back(SubData);
         }

clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp

void foo(double *&arg) {
     }
     // CK2:     call void @__tgt_target_data_end{{.+}}[[MTYPE00]]
     // CK2:     [[DECL:%.+]] = getelementptr inbounds [[ST]], [[ST]]* %this1, i32 0, i32 0
     // CK2:     [[TTT:%.+]] = load double*, double** [[DECL]],
     // CK2:     getelementptr inbounds double, double* [[TTT]], i32 1
     a++;
 
     // CK2:     [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 1
-    // CK2:     [[CBP:%.+]] = bitcast i8** [[BP]] to double***
-    // CK2:     store double** [[RVAL:%.+]], double*** [[CBP]],
+    // CK2:     [[CBP:%.+]] = bitcast i8** [[BP]] to double****
+    // CK2:     store double*** [[RVAL:%.+]], double**** [[CBP]],
     // CK2:     call void @__tgt_target_data_begin{{.+}}[[MTYPE01]]
-    // CK2:     [[CBP1:%.+]] = bitcast double*** [[CBP]] to double**
+    // CK2:     [[CBP1:%.+]] = bitcast double**** [[CBP]] to double**
     // CK2:     [[VAL:%.+]] = load double*, double** [[CBP1]],
     // CK2:     store double* [[VAL]], double** [[PVT:%.+]],
     // CK2:     store double** [[PVT]], double*** [[PVT2:%.+]],
     // CK2:     [[TT1:%.+]] = load double**, double*** [[PVT2]],
     // CK2:     [[TT2:%.+]] = load double*, double** [[TT1]],
     // CK2:     getelementptr inbounds double, double* [[TT2]], i32 1
     #pragma omp target data map(b[:10]) use_device_ptr(b)
     {
void foo(double *&arg) {
     // CK2:     call void @__tgt_target_data_end{{.+}}[[MTYPE02]]
     // CK2:     [[DECL:%.+]] = getelementptr inbounds [[ST]], [[ST]]* %this1, i32 0, i32 0
     // CK2:     [[TTT:%.+]] = load double*, double** [[DECL]],
     // CK2:     getelementptr inbounds double, double* [[TTT]], i32 1
     a++;
     la++;
 
     // CK2:     [[BP1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %{{.+}}, i32 0, i32 1
-    // CK2:     [[CBP1:%.+]] = bitcast i8** [[BP1]] to double***
-    // CK2:     store double** [[RVAL1:%.+]], double*** [[CBP1]],
+    // CK2:     [[CBP1:%.+]] = bitcast i8** [[BP1]] to double****
+    // CK2:     store double*** [[RVAL1:%.+]], double**** [[CBP1]],
     // CK2:     [[BP2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %{{.+}}, i32 0, i32 2
     // CK2:     [[CBP2:%.+]] = bitcast i8** [[BP2]] to double***
     // CK2:     store double** [[RVAL2:%.+]], double*** [[CBP2]],
     // CK2:     call void @__tgt_target_data_begin{{.+}}[[MTYPE03]]
     // CK2:     [[_CBP2:%.+]] = bitcast double*** [[CBP2]] to double**
     // CK2:     [[VAL2:%.+]] = load double*, double** [[_CBP2]],
     // CK2:     store double* [[VAL2]], double** [[PVT2:%.+]],
     // CK2:     store double** [[PVT2]], double*** [[_PVT2:%.+]],
-    // CK2:     [[_CBP1:%.+]] = bitcast double*** [[CBP1]] to double**
+    // CK2:     [[_CBP1:%.+]] = bitcast double**** [[CBP1]] to double**
     // CK2:     [[VAL1:%.+]] = load double*, double** [[_CBP1]],
     // CK2:     store double* [[VAL1]], double** [[PVT1:%.+]],
     // CK2:     store double** [[PVT1]], double*** [[_PVT1:%.+]],
     // CK2:     [[TT2:%.+]] = load double**, double*** [[_PVT2]],
     // CK2:     [[_TT2:%.+]] = load double*, double** [[TT2]],
     // CK2:     getelementptr inbounds double, double* [[_TT2]], i32 1
     // CK2:     [[TT1:%.+]] = load double**, double*** [[_PVT1]],
     // CK2:     [[_TT1:%.+]] = load double*, double** [[TT1]],

clang/test/OpenMP/target_map_codegen_28.cpp

 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY28 %s
 // SIMD-ONLY28-NOT: {{__kmpc|__tgt}}
 #ifdef CK29
 
 // CK29: [[SSA:%.+]] = type { double*, double** }
 // CK29: [[SSB:%.+]]  = type { [[SSA]]*, [[SSA]]** }
 
 // CK29-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
-// CK29: [[MTYPE00:@.+]] = private {{.*}}constant [3 x i64] [i64 32, i64 281474976710672, i64 19]
+// CK29: [[MTYPE00:@.+]] = private {{.*}}constant [2 x i64] [i64 32, i64 281474976710675]
 
 // CK29-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK29: [[MTYPE01:@.+]] = private {{.*}}constant [3 x i64] [i64 32, i64 281474976710672, i64 19]
 
 // CK29-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
-// CK29: [[MTYPE02:@.+]] = private {{.*}}constant [3 x i64] [i64 32, i64 281474976710672, i64 19]
+// CK29: [[MTYPE02:@.+]] = private {{.*}}constant [2 x i64] [i64 32, i64 281474976710675]
 
 struct SSA{
   double *p;
   double *≺
   SSA(double *&pr) : pr(pr) {}
 };
 
 struct SSB{
   SSA *p;
   SSA *≺
   SSB(SSA *&pr) : pr(pr) {}
 
   // CK29-LABEL: define {{.+}}foo
   void foo() {
 
     // Region 00
-    // CK29-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{[^,]+}}, i8* {{[^,]+}}, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z:64|32]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null, i8** null)
+    // CK29-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z:64|32]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null, i8** null)
 
     // CK29-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
     // CK29-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
     // CK29-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 
     // CK29-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
     // CK29-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
     // CK29-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
     // CK29-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[SSB]]**
-    // CK29-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SSA]]**
+    // CK29-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SSA]]***
     // CK29-DAG: store [[SSB]]* [[VAR0:%.+]], [[SSB]]** [[CBP0]]
     // CK29-DAG: store [[SSA]]** [[VAR00:%.+]], [[SSA]]*** [[CP0]]
     // CK29-DAG: store i64 %{{.+}}, i64* [[S0]]
     // CK29-DAG: [[VAR0]] = load [[SSB]]*, [[SSB]]** %
     // CK29-DAG: [[VAR00]] = getelementptr inbounds [[SSB]], [[SSB]]* [[VAR0]], i32 0, i32 0
 
-    // CK29-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
-    // CK29-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
-    // CK29-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
-    // CK29-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[SSA]]***
-    // CK29-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double***
-    // CK29-DAG: store [[SSA]]** [[VAR00]], [[SSA]]*** [[CBP1]]
-    // CK29-DAG: store double** [[VAR1:%.+]], double*** [[CP1]]
-    // CK29-DAG: store i64 {{8|4}}, i64* [[S1]]
-    // CK29-DAG: [[VAR1]] = load double**, double*** [[VAR1_REF:%.+]],
-    // CK29-DAG: [[VAR1_REF]] = getelementptr inbounds [[SSA]], [[SSA]]* %{{.+}}, i32 0, i32 1
-
-    // CK29-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
-    // CK29-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
-    // CK29-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2
-    // CK29-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to double***
+    // CK29-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+    // CK29-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+    // CK29-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
+    // CK29-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to double****
     // CK29-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to double**
-    // CK29-DAG: store double** [[VAR1]], double*** [[CBP2]]
+    // CK29-DAG: store double*** [[VAR1:%.+]], double**** [[CBP2]]
     // CK29-DAG: store double* [[VAR2:%.+]], double** [[CP2]]
     // CK29-DAG: store i64 80, i64* [[S2]]
+    // CK29-DAG: [[VAR1]] = getelementptr inbounds [[SSA]], [[SSA]]* %{{.+}}, i32 0, i32 1
     // CK29-DAG: [[VAR2]] = getelementptr inbounds double, double* [[VAR22:%.+]], i{{.+}} 0
     // CK29-DAG: [[VAR22]] = load double*, double** %{{.+}},
 
     // CK29: call void [[CALL00:@.+]]([[SSB]]* {{[^,]+}})
     #pragma omp target map(p->pr[:10])
     {
       p->pr++;
     }
 
     // Region 01
     // CK29-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{[^,]+}}, i8* {{[^,]+}}, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE01]]{{.+}}, i8** null, i8** null)
 
     // CK29-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
     // CK29-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
     // CK29-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 
     // CK29-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
     // CK29-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
     // CK29-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
     // CK29-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[SSB]]**
-    // CK29-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SSA]]**
+    // CK29-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SSA]]****
     // CK29-DAG: store [[SSB]]* [[VAR0]], [[SSB]]** [[CBP0]]
-    // CK29-DAG: store [[SSA]]** [[VAR00:%.+]], [[SSA]]*** [[CP0]]
+    // CK29-DAG: store [[SSA]]*** [[VAR000:%.+]], [[SSA]]**** [[CP0]]
     // CK29-DAG: store i64 %{{.+}}, i64* [[S0]]
-    // CK29-DAG: [[VAR00]] = load [[SSA]]**, [[SSA]]*** [[VAR000:%.+]],
     // CK29-DAG: [[VAR000]] = getelementptr inbounds [[SSB]], [[SSB]]* [[VAR0]], i32 0, i32 1
 
     // CK29-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
     // CK29-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
     // CK29-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
-    // CK29-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[SSA]]***
+    // CK29-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[SSA]]****
     // CK29-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double***
-    // CK29-DAG: store [[SSA]]** [[VAR00]], [[SSA]]*** [[CBP1]]
+    // CK29-DAG: store [[SSA]]*** [[VAR000]], [[SSA]]**** [[CBP1]]
     // CK29-DAG: store double** [[VAR1:%.+]], double*** [[CP1]]
     // CK29-DAG: store i64 {{8|4}}, i64* [[S1]]
     // CK29-DAG: [[VAR1]] = getelementptr inbounds [[SSA]], [[SSA]]* %{{.+}}, i32 0, i32 0
 
     // CK29-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
     // CK29-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
     // CK29-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2
     // CK29-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to double***
     // CK29-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to double**
     // CK29-DAG: store double** [[VAR1]], double*** [[CBP2]]
     // CK29-DAG: store double* [[VAR2:%.+]], double** [[CP2]]
     // CK29-DAG: store i64 80, i64* [[S2]]
     // CK29-DAG: [[VAR2]] = getelementptr inbounds double, double* [[VAR22:%.+]], i{{.+}} 0
     // CK29-DAG: [[VAR22]] = load double*, double** %{{.+}},
 
     // CK29: call void [[CALL00:@.+]]([[SSB]]* {{[^,]+}})
     #pragma omp target map(pr->p[:10])
     {
       pr->p++;
     }
 
     // Region 02
-    // CK29-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{[^,]+}}, i8* {{[^,]+}}, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE02]]{{.+}}, i8** null, i8** null)
+    // CK29-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE02]]{{.+}}, i8** null, i8** null)
 
     // CK29-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
     // CK29-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
     // CK29-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 
     // CK29-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
     // CK29-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
     // CK29-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
     // CK29-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[SSB]]**
-    // CK29-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SSA]]**
+    // CK29-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SSA]]****
     // CK29-DAG: store [[SSB]]* [[VAR0]], [[SSB]]** [[CBP0]]
-    // CK29-DAG: store [[SSA]]** [[VAR00:%.+]], [[SSA]]*** [[CP0]]
+    // CK29-DAG: store [[SSA]]*** [[VAR000:%.+]], [[SSA]]**** [[CP0]]
     // CK29-DAG: store i64 %{{.+}}, i64* [[S0]]
-    // CK29-DAG: [[VAR00]] = load [[SSA]]**, [[SSA]]*** [[VAR000:%.+]],
     // CK29-DAG: [[VAR000]] = getelementptr inbounds [[SSB]], [[SSB]]* [[VAR0]], i32 0, i32 1
 
-    // CK29-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
-    // CK29-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
-    // CK29-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
-    // CK29-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[SSA]]***
-    // CK29-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double***
-    // CK29-DAG: store [[SSA]]** [[VAR00]], [[SSA]]*** [[CBP1]]
-    // CK29-DAG: store double** [[VAR1:%.+]], double*** [[CP1]]
-    // CK29-DAG: store i64 {{8|4}}, i64* [[S1]]
-    // CK29-DAG: [[VAR1]] = load double**, double*** [[VAR1_REF:%.+]],
-    // CK29-DAG: [[VAR1_REF]] = getelementptr inbounds [[SSA]], [[SSA]]* %{{.+}}, i32 0, i32 1
-
-    // CK29-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
-    // CK29-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
-    // CK29-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2
-    // CK29-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to double***
+    // CK29-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+    // CK29-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+    // CK29-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
+    // CK29-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to double****
     // CK29-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to double**
-    // CK29-DAG: store double** [[VAR1]], double*** [[CBP2]]
+    // CK29-DAG: store double*** [[VAR1:%.+]], double**** [[CBP2]]
     // CK29-DAG: store double* [[VAR2:%.+]], double** [[CP2]]
     // CK29-DAG: store i64 80, i64* [[S2]]
+    // CK29-DAG: [[VAR1]] = getelementptr inbounds [[SSA]], [[SSA]]* %{{.+}}, i32 0, i32 1
     // CK29-DAG: [[VAR2]] = getelementptr inbounds double, double* [[VAR22:%.+]], i{{.+}} 0
     // CK29-DAG: [[VAR22]] = load double*, double** %{{.+}},
 
     // CK29: call void [[CALL00:@.+]]([[SSB]]* {{[^,]+}})
     #pragma omp target map(pr->pr[:10])
     {
       pr->pr++;
     }

clang/test/OpenMP/target_map_codegen_35.cpp

+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+///==========================================================================///
+// RUN: %clang_cc1 -DCK35 -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 CK35 --check-prefix CK35-64
+// RUN: %clang_cc1 -DCK35 -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 CK35 --check-prefix CK35-64
+// RUN: %clang_cc1 -DCK35 -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 CK35 --check-prefix CK35-32
+// RUN: %clang_cc1 -DCK35 -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 CK35 --check-prefix CK35-32
+
+// RUN: %clang_cc1 -DCK35 -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 -DCK35 -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 -DCK35 -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 -DCK35 -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 CK35
+
+class S {
+public:
+  S(double &b) : b(b) {}
+  int a;
+  double &b;
+  void foo();
+};
+
+// TARGET_PARAM = 0x20
+// MEMBER_OF_1 | TO = 0x1000000000001
+// MEMBER_OF_1 | PTR_AND_OBJ | TO = 0x1000000000011
+// CK35-DAG: [[MTYPE_TO:@.+]] = {{.+}}constant [4 x i64] [i64 [[#0x20]], i64 [[#0x1000000000001]], i64 [[#0x1000000000001]], i64 [[#0x1000000000011]]]
+// TARGET_PARAM = 0x20
+// MEMBER_OF_1 | PTR_AND_OBJ | FROM = 0x1000000000012
+// CK35-DAG: [[MTYPE_FROM:@.+]] = {{.+}}constant [2 x i64] [i64 [[#0x20]], i64 [[#0x1000000000012]]]
+
+void ref_map() {
+  double b;
+  S s(b);
+
+  // CK35-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** null)
+  // CK35-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK35-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+  // CK35-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+
+  // pass TARGET_PARAM {&s, &s, ((void*)(&s+1)-(void*)&s)}
+
+  // CK35-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK35-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK35-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+
+  // CK35-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to %class.S**
+  // CK35-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to %class.S**
+
+  // CK35-DAG: store %class.S* [[S_ADDR:%.+]], %class.S** [[BPC0]],
+  // CK35-DAG: store %class.S* [[S_ADDR]], %class.S** [[PC0]],
+  // CK35-DAG: store i64 [[S_SIZE:%.+]], i64* [[S0]],
+
+  // CK35-DAG: [[S_SIZE]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+  // CK35-DAG: [[SZ]] = sub i64 [[S_1_INTPTR:%.+]], [[S_INTPTR:%.+]]
+  // CK35-DAG: [[S_1_INTPTR]] = ptrtoint i8* [[S_1_VOID:%.+]] to i64
+  // CK35-DAG: [[S_INTPTR]] = ptrtoint i8* [[S_VOID:%.+]] to i64
+  // CK35-DAG: [[S_1_VOID]] = bitcast %class.S* [[S_1:%.+]] to i8*
+  // CK35-DAG: [[S_VOID]] = bitcast %class.S* [[S_ADDR]] to i8*
+  // CK35-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.
+
+  // CK35-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+  // CK35-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK35-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
+
+  // CK35-DAG: [[BPC1:%.+]] = bitcast i8** [[BP1]] to %class.S**
+  // CK35-DAG: [[PC1:%.+]] = bitcast i8** [[P1]] to %class.S**
+
+  // CK35-DAG: store %class.S* [[S_ADDR]], %class.S** [[BPC1]],
+  // CK35-DAG: store %class.S* [[S_ADDR]], %class.S** [[PC1]],
+  // CK35-DAG: store i64 [[A_SIZE:%.+]], i64* [[S1]],
+
+  // CK35-DAG: [[A_SIZE]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+  // CK35-DAG: [[SZ]] = sub i64 [[B_BEGIN_INTPTR:%.+]], [[S_INTPTR:%.+]]
+  // CK35-DAG: [[S_INTPTR]] = ptrtoint i8* [[S_VOID:%.+]] to i64
+  // CK35-DAG: [[B_BEGIN_INTPTR]] = ptrtoint i8* [[B_BEGIN_VOID:%.+]] to i64
+  // CK35-DAG: [[S_VOID]] = bitcast %class.S* [[S_ADDR]] to i8*
+  // CK35-DAG: [[B_BEGIN_VOID]] = bitcast double** [[B_ADDR:%.+]] to i8*
+  // CK35-DAG: [[B_ADDR]] = getelementptr inbounds %class.S, %class.S* [[S_ADDR]], i32 0, i32 1
+
+  // pass MEMBER_OF_1 | TO {&s, &s.b+1, ((void*)(&s+1)-(void*)(&s.b+1))} to copy the data of remainder of s.
+
+  // CK35-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
+  // CK35-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
+  // CK35-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2
+
+  // CK35-DAG: [[BPC2:%.+]] = bitcast i8** [[BP2]] to %class.S**
+  // CK35-DAG: [[PC2:%.+]] = bitcast i8** [[P2]] to double***
+
+  // CK35-DAG: store %class.S* [[S_ADDR]], %class.S** [[BPC2]],
+  // CK35-DAG: store double** [[B_END:%.+]], double*** [[PC2]],
+  // CK35-DAG: store i64 [[REM_SIZE:%.+]], i64* [[S2]],
+
+  // CK35-DAG: [[B_END]] = getelementptr double*, double** [[B_ADDR]], i{{.+}} 1
+
+  // CK35-DAG: [[REM_SIZE]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+  // CK35-DAG: [[SZ]] = sub i64 [[S_END_INTPTR:%.+]], [[B_END_INTPTR:%.+]]
+  // CK35-DAG: [[B_END_INTPTR]] = ptrtoint i8* [[B_END_VOID:%.+]] to i64
+  // CK35-DAG: [[S_END_INTPTR]] = ptrtoint i8* [[S_END_VOID:%.+]] to i64
+  // CK35-DAG: [[B_END_VOID]] = bitcast double** [[B_END]] to i8*
+  // CK35-DAG: [[S_END_VOID]] = getelementptr i8, i8* [[S_LAST:%.+]], i{{.+}} 1
+  // CK35-64-DAG: [[S_LAST]] = getelementptr i8, i8* [[S_VOIDPTR:%.+]], i64 15
+  // CK35-32-DAG: [[S_LAST]] = getelementptr i8, i8* [[S_VOIDPTR:%.+]], i32 7
+  // CK35-DAG: [[S_VOIDPTR]] = bitcast %class.S* [[S_ADDR]] to i8*
+
+  // pass MEMBER_OF_1 | PTR_AND_OBJ | TO {&s, &s.b, 8|4} to copy the data of s.b.
+
+  // CK35-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3
+  // CK35-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3
+  // CK35-DAG: [[S3:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 3
+
+  // CK35-DAG: [[BPC3:%.+]] = bitcast i8** [[BP3]] to %class.S**
+  // CK35-DAG: [[PC3:%.+]] = bitcast i8** [[P3]] to double**
+
+  // CK35-DAG: store %class.S* [[S_ADDR]], %class.S** [[BPC3]],
+  // CK35-DAG: store double* [[B_ADDR:%.+]], double** [[PC3]],
+  // CK35-DAG: store i64 8, i64* [[S3]],
+
+  // CK35-DAG: [[B_ADDR]] = load double*, double** [[B_REF:%.+]],
+  // CK35-DAG: [[B_REF]] = getelementptr inbounds %class.S, %class.S* [[S_ADDR]], i32 0, i32 1
+
+  #pragma omp target map(to: s, s.b)
+  s.foo();
+
+  // CK35 : call void
+
+  // CK35-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 -1, i8* @{{.+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE_FROM]]{{.+}}, i8** null, i8** null)
+  // CK35-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK35-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+  // CK35-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+
+  // pass TARGET_PARAM {&s, &s.b, ((void*)(&s.b+1)-(void*)&s.b)}
+
+  // CK35-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK35-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK35-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+
+  // CK35-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to %class.S**
+  // CK35-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to double***
+
+  // CK35-DAG: store %class.S* [[S_ADDR]], %class.S** [[BPC0]],
+  // CK35-DAG: store double** [[SB_ADDR:%.+]], double*** [[PC0]],
+  // CK35-DAG: store i64 [[B_SIZE:%.+]], i64* [[S0]],
+
+  // CK35-DAG: [[B_SIZE]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+  // CK35-DAG: [[SZ]] = sub i64 [[SB_1_INTPTR:%.+]], [[SB_INTPTR:%.+]]
+  // CK35-DAG: [[SB_1_INTPTR]] = ptrtoint i8* [[SB_1_VOID:%.+]] to i64
+  // CK35-DAG: [[SB_INTPTR]] = ptrtoint i8* [[SB_VOID:%.+]] to i64
+  // CK35-DAG: [[SB_1_VOID]] = bitcast double** [[SB_1:%.+]] to i8*
+  // CK35-DAG: [[SB_VOID]] = bitcast double** [[SB_ADDR:%.+]] to i8*
+  // CK35-DAG: [[SB_ADDR]] = getelementptr inbounds %class.S, %class.S* [[S_ADDR]], i32 0, i32 1
+  // CK35-DAG: [[SB_1]] = getelementptr double*, double** [[SB_ADDR]], i{{.+}} 1
+
+  // pass MEMBER_OF_1 | PTR_AND_OBJ | FROM {&s, &s.b, 8|4} to copy the data of s.c.
+
+  // CK35-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+  // CK35-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK35-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
+
+  // CK35-DAG: [[BPC1:%.+]] = bitcast i8** [[BP1]] to %class.S**
+  // CK35-DAG: [[PC1:%.+]] = bitcast i8** [[P1]] to double**
+
+  // CK35-DAG: store %class.S* [[S_ADDR]], %class.S** [[BPC1]],
+  // CK35-DAG: store double* [[B_ADDR:%.+]], double** [[PC1]],
+  // CK35-DAG: store i64 8, i64* [[S1]],
+
+  // CK35-DAG: [[B_ADDR]] = load double*, double** [[SB_ADDR]],
+
+  #pragma omp target map(from: s.b)
+  s.foo();
+}
+
+#endif // CK35
+#endif

openmp/libomptarget/test/mapping/data_member_ref.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 <stdio.h>
+
+struct View {
+  int Data;
+};
+
+struct ViewPtr {
+  int *Data;
+};
+
+template <typename T> struct Foo {
+  Foo(T &V) : VRef(V) {}
+  T &VRef;
+};
+
+int main() {
+  View V;
+  V.Data = 123456;
+  Foo<View> Bar(V);
+  ViewPtr V1;
+  int Data = 123456;
+  V1.Data = &Data;
+  Foo<ViewPtr> Baz(V1);
+
+  // CHECK: Host 123456.
+  printf("Host %d.\n", Bar.VRef.Data);
+#pragma omp target map(Bar.VRef)
+  {
+    // CHECK: Device 123456.
+    printf("Device %d.\n", Bar.VRef.Data);
+    V.Data = 654321;
+    // CHECK: Device 654321.
+    printf("Device %d.\n", Bar.VRef.Data);
+  }
+  // CHECK: Host 654321 654321.
+  printf("Host %d %d.\n", Bar.VRef.Data, V.Data);
+  V.Data = 123456;
+  // CHECK: Host 123456.
+  printf("Host %d.\n", Bar.VRef.Data);
+#pragma omp target map(Bar) map(Bar.VRef)
+  {
+    // CHECK: Device 123456.
+    printf("Device %d.\n", Bar.VRef.Data);
+    V.Data = 654321;
+    // CHECK: Device 654321.
+    printf("Device %d.\n", Bar.VRef.Data);
+  }
+  // CHECK: Host 654321 654321.
+  printf("Host %d %d.\n", Bar.VRef.Data, V.Data);
+  // CHECK: Host 123456.
+  printf("Host %d.\n", *Baz.VRef.Data);
+#pragma omp target map(*Baz.VRef.Data)
+  {
+    // CHECK: Device 123456.
+    printf("Device %d.\n", *Baz.VRef.Data);
+    *V1.Data = 654321;
+    // CHECK: Device 654321.
+    printf("Device %d.\n", *Baz.VRef.Data);
+  }
+  // CHECK: Host 654321 654321 654321.
+  printf("Host %d %d %d.\n", *Baz.VRef.Data, *V1.Data, Data);
+  return 0;
+}