[OPENMP] Add support for mapping memory pointed by member pointer.

Added support for map(s, s.ptr[0:1]) kind of mapping.

llvm-svn: 342648

clang/lib/CodeGen/CGOpenMPRuntime.cpp

@@ -6752,7 +6752,9 @@ private:
       MapBaseValuesArrayTy &BasePointers, MapValuesArrayTy &Pointers,
       MapValuesArrayTy &Sizes, MapFlagsArrayTy &Types,
       StructRangeInfoTy &PartialStruct, bool IsFirstComponentList,
-      bool IsImplicit) const {
+      bool IsImplicit,
+      ArrayRef<OMPClauseMappableExprCommon::MappableExprComponentListRef>
+          OverlappedElements = llvm::None) const {
     // The following summarizes what has to be generated for each map and the
     // types below. The generated information is expressed in this order:
     // base pointer, section pointer, size, flags
@@ -7023,7 +7025,6 @@ private:
 
         Address LB =
             CGF.EmitOMPSharedLValue(I->getAssociatedExpression()).getAddress();
-        llvm::Value *Size = getExprTypeSize(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
@@ -7032,6 +7033,70 @@ private:
             IsPointer && EncounteredME &&
             (dyn_cast<MemberExpr>(I->getAssociatedExpression()) ==
              EncounteredME);
+        if (!OverlappedElements.empty()) {
+          // 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, CharUnits::One());
+          PartialStruct.HighestElem = {
+              std::numeric_limits<decltype(
+                  PartialStruct.HighestElem.first)>::max(),
+              HB};
+          PartialStruct.Base = BP;
+          // Emit data for non-overlapped data.
+          OpenMPOffloadMappingFlags Flags =
+              OMP_MAP_MEMBER_OF |
+              getMapTypeBits(MapType, MapTypeModifier, IsImplicit,
+                             /*AddPtrFlag=*/false,
+                             /*AddIsTargetParamFlag=*/false);
+          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()) {
+                ComponentLB =
+                    CGF.EmitOMPSharedLValue(MC.getAssociatedExpression())
+                        .getAddress();
+                Size = CGF.Builder.CreatePtrDiff(
+                    CGF.EmitCastToVoidPtr(ComponentLB.getPointer()),
+                    CGF.EmitCastToVoidPtr(LB.getPointer()));
+                break;
+              }
+            }
+            BasePointers.push_back(BP.getPointer());
+            Pointers.push_back(LB.getPointer());
+            Sizes.push_back(Size);
+            Types.push_back(Flags);
+            LB = CGF.Builder.CreateConstGEP(ComponentLB, 1,
+                                            CGF.getPointerSize());
+          }
+          BasePointers.push_back(BP.getPointer());
+          Pointers.push_back(LB.getPointer());
+          Size = CGF.Builder.CreatePtrDiff(
+              CGF.EmitCastToVoidPtr(
+                  CGF.Builder.CreateConstGEP(HB, 1, CharUnits::One())
+                      .getPointer()),
+              CGF.EmitCastToVoidPtr(LB.getPointer()));
+          Sizes.push_back(Size);
+          Types.push_back(Flags);
+          break;
+        }
+        llvm::Value *Size = getExprTypeSize(I->getAssociatedExpression());
         if (!IsMemberPointer) {
           BasePointers.push_back(BP.getPointer());
           Pointers.push_back(LB.getPointer());
@@ -7136,6 +7201,66 @@ private:
     Flags |= MemberOfFlag;
   }
 
+  void getPlainLayout(const CXXRecordDecl *RD,
+                      llvm::SmallVectorImpl<const FieldDecl *> &Layout,
+                      bool AsBase) const {
+    const CGRecordLayout &RL = CGF.getTypes().getCGRecordLayout(RD);
+
+    llvm::StructType *St =
+        AsBase ? RL.getBaseSubobjectLLVMType() : RL.getLLVMType();
+
+    unsigned NumElements = St->getNumElements();
+    llvm::SmallVector<
+        llvm::PointerUnion<const CXXRecordDecl *, const FieldDecl *>, 4>
+        RecordLayout(NumElements);
+
+    // Fill bases.
+    for (const auto &I : RD->bases()) {
+      if (I.isVirtual())
+        continue;
+      const auto *Base = I.getType()->getAsCXXRecordDecl();
+      // Ignore empty bases.
+      if (Base->isEmpty() || CGF.getContext()
+                                 .getASTRecordLayout(Base)
+                                 .getNonVirtualSize()
+                                 .isZero())
+        continue;
+
+      unsigned FieldIndex = RL.getNonVirtualBaseLLVMFieldNo(Base);
+      RecordLayout[FieldIndex] = Base;
+    }
+    // Fill in virtual bases.
+    for (const auto &I : RD->vbases()) {
+      const auto *Base = I.getType()->getAsCXXRecordDecl();
+      // Ignore empty bases.
+      if (Base->isEmpty())
+        continue;
+      unsigned FieldIndex = RL.getVirtualBaseIndex(Base);
+      if (RecordLayout[FieldIndex])
+        continue;
+      RecordLayout[FieldIndex] = Base;
+    }
+    // Fill in all the fields.
+    assert(!RD->isUnion() && "Unexpected union.");
+    for (const auto *Field : RD->fields()) {
+      // Fill in non-bitfields. (Bitfields always use a zero pattern, which we
+      // will fill in later.)
+      if (!Field->isBitField()) {
+        unsigned FieldIndex = RL.getLLVMFieldNo(Field);
+        RecordLayout[FieldIndex] = Field;
+      }
+    }
+    for (const llvm::PointerUnion<const CXXRecordDecl *, const FieldDecl *>
+             &Data : RecordLayout) {
+      if (Data.isNull())
+        continue;
+      if (const auto *Base = Data.dyn_cast<const CXXRecordDecl *>())
+        getPlainLayout(Base, Layout, /*AsBase=*/true);
+      else
+        Layout.push_back(Data.get<const FieldDecl *>());
+    }
+  }
+
 public:
   MappableExprsHandler(const OMPExecutableDirective &Dir, CodeGenFunction &CGF)
       : CurDir(Dir), CGF(CGF) {
@@ -7376,9 +7501,6 @@ public:
            "Not expecting to generate map info for a variable array type!");
 
     // We need to know when we generating information for the first component
-    // associated with a capture, because the mapping flags depend on it.
-    bool IsFirstComponentList = true;
-
     const ValueDecl *VD = Cap->capturesThis()
                               ? nullptr
                               : Cap->getCapturedVar()->getCanonicalDecl();
@@ -7394,17 +7516,143 @@ public:
       return;
     }
 
+    using MapData =
+        std::tuple<OMPClauseMappableExprCommon::MappableExprComponentListRef,
+                   OpenMPMapClauseKind, OpenMPMapClauseKind, bool>;
+    SmallVector<MapData, 4> DeclComponentLists;
     // FIXME: MSVC 2013 seems to require this-> to find member CurDir.
-    for (const auto *C : this->CurDir.getClausesOfKind<OMPMapClause>())
+    for (const auto *C : this->CurDir.getClausesOfKind<OMPMapClause>()) {
       for (const auto &L : C->decl_component_lists(VD)) {
         assert(L.first == VD &&
                "We got information for the wrong declaration??");
         assert(!L.second.empty() &&
                "Not expecting declaration with no component lists.");
-        generateInfoForComponentList(C->getMapType(), C->getMapTypeModifier(),
+        DeclComponentLists.emplace_back(L.second, C->getMapType(),
-                                     L.second, BasePointers, Pointers, Sizes,
+                                        C->getMapTypeModifier(),
-                                     Types, PartialStruct, IsFirstComponentList,
                                         C->isImplicit());
+      }
+    }
+
+    // Find overlapping elements (including the offset from the base element).
+    llvm::SmallDenseMap<
+        const MapData *,
+        llvm::SmallVector<
+            OMPClauseMappableExprCommon::MappableExprComponentListRef, 4>,
+        4>
+        OverlappedData;
+    size_t Count = 0;
+    for (const MapData &L : DeclComponentLists) {
+      OMPClauseMappableExprCommon::MappableExprComponentListRef Components;
+      OpenMPMapClauseKind MapType;
+      OpenMPMapClauseKind MapTypeModifier;
+      bool IsImplicit;
+      std::tie(Components, MapType, MapTypeModifier, IsImplicit) = L;
+      ++Count;
+      for (const MapData &L1 : makeArrayRef(DeclComponentLists).slice(Count)) {
+        OMPClauseMappableExprCommon::MappableExprComponentListRef Components1;
+        std::tie(Components1, MapType, MapTypeModifier, IsImplicit) = L1;
+        auto CI = Components.rbegin();
+        auto CE = Components.rend();
+        auto SI = Components1.rbegin();
+        auto SE = Components1.rend();
+        for (; CI != CE && SI != SE; ++CI, ++SI) {
+          if (CI->getAssociatedExpression()->getStmtClass() !=
+              SI->getAssociatedExpression()->getStmtClass())
+            break;
+          // Are we dealing with different variables/fields?
+          if (CI->getAssociatedDeclaration() != SI->getAssociatedDeclaration())
+            break;
+        }
+        // Found overlapping if, at least for one component, reached the head of
+        // the components list.
+        if (CI == CE || SI == SE) {
+          assert((CI != CE || SI != SE) &&
+                 "Unexpected full match of the mapping components.");
+          const MapData &BaseData = CI == CE ? L : L1;
+          OMPClauseMappableExprCommon::MappableExprComponentListRef SubData =
+              SI == SE ? Components : Components1;
+          auto It = CI == CE ? SI : CI;
+          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())
+        getPlainLayout(CRD, Layout, /*AsBase=*/false);
+      else {
+        const auto *RD = VD->getType().getCanonicalType()->getAsRecordDecl();
+        Layout.append(RD->field_begin(), RD->field_end());
+      }
+    }
+    for (auto &Pair : OverlappedData) {
+      llvm::sort(
+          Pair.getSecond(),
+          [&Layout](
+              OMPClauseMappableExprCommon::MappableExprComponentListRef First,
+              OMPClauseMappableExprCommon::MappableExprComponentListRef
+                  Second) {
+            auto CI = First.rbegin();
+            auto CE = First.rend();
+            auto SI = Second.rbegin();
+            auto SE = Second.rend();
+            for (; CI != CE && SI != SE; ++CI, ++SI) {
+              if (CI->getAssociatedExpression()->getStmtClass() !=
+                  SI->getAssociatedExpression()->getStmtClass())
+                break;
+              // Are we dealing with different variables/fields?
+              if (CI->getAssociatedDeclaration() !=
+                  SI->getAssociatedDeclaration())
+                break;
+            }
+            assert(CI != CE && SI != SE &&
+                   "Unexpected end of the map components.");
+            const auto *FD1 = cast<FieldDecl>(CI->getAssociatedDeclaration());
+            const auto *FD2 = cast<FieldDecl>(SI->getAssociatedDeclaration());
+            if (FD1->getParent() == FD2->getParent())
+              return FD1->getFieldIndex() < FD2->getFieldIndex();
+            const auto It =
+                llvm::find_if(Layout, [FD1, FD2](const FieldDecl *FD) {
+                  return FD == FD1 || FD == FD2;
+                });
+            return *It == FD1;
+          });
+    }
+
+    // Associated with a capture, because the mapping flags depend on it.
+    // Go through all of the elements with the overlapped elements.
+    for (const auto &Pair : OverlappedData) {
+      const MapData &L = *Pair.getFirst();
+      OMPClauseMappableExprCommon::MappableExprComponentListRef Components;
+      OpenMPMapClauseKind MapType;
+      OpenMPMapClauseKind MapTypeModifier;
+      bool IsImplicit;
+      std::tie(Components, MapType, MapTypeModifier, IsImplicit) = L;
+      ArrayRef<OMPClauseMappableExprCommon::MappableExprComponentListRef>
+          OverlappedComponents = Pair.getSecond();
+      bool IsFirstComponentList = true;
+      generateInfoForComponentList(MapType, MapTypeModifier, Components,
+                                   BasePointers, Pointers, Sizes, Types,
+                                   PartialStruct, IsFirstComponentList,
+                                   IsImplicit, OverlappedComponents);
+    }
+    // Go through other elements without overlapped elements.
+    bool IsFirstComponentList = OverlappedData.empty();
+    for (const MapData &L : DeclComponentLists) {
+      OMPClauseMappableExprCommon::MappableExprComponentListRef Components;
+      OpenMPMapClauseKind MapType;
+      OpenMPMapClauseKind MapTypeModifier;
+      bool IsImplicit;
+      std::tie(Components, MapType, MapTypeModifier, IsImplicit) = L;
+      auto It = OverlappedData.find(&L);
+      if (It == OverlappedData.end())
+        generateInfoForComponentList(MapType, MapTypeModifier, Components,
+                                     BasePointers, Pointers, Sizes, Types,
+                                     PartialStruct, IsFirstComponentList,
+                                     IsImplicit);
       IsFirstComponentList = false;
     }
   }

clang/lib/Sema/SemaOpenMP.cpp

@@ -12333,6 +12333,26 @@ static bool checkMapConflicts(
         // An expression is a subset of the other.
         if (CurrentRegionOnly && (CI == CE || SI == SE)) {
           if (CKind == OMPC_map) {
+            if (CI != CE || SI != SE) {
+              // Allow constructs like this: map(s, s.ptr[0:1]), where s.ptr is
+              // a pointer.
+              auto Begin =
+                  CI != CE ? CurComponents.begin() : StackComponents.begin();
+              auto End = CI != CE ? CurComponents.end() : StackComponents.end();
+              auto It = Begin;
+              while (It != End && !It->getAssociatedDeclaration())
+                std::advance(It, 1);
+              assert(It != End &&
+                     "Expected at least one component with the declaration.");
+              if (It != Begin && It->getAssociatedDeclaration()
+                                     ->getType()
+                                     .getCanonicalType()
+                                     ->isAnyPointerType()) {
+                IsEnclosedByDataEnvironmentExpr = false;
+                EnclosingExpr = nullptr;
+                return false;
+              }
+            }
             SemaRef.Diag(ELoc, diag::err_omp_map_shared_storage) << ERange;
           } else {
             assert(CKind == OMPC_to || CKind == OMPC_from);

clang/test/OpenMP/target_map_codegen.cpp

@@ -5100,5 +5100,162 @@ void explicit_maps_member_pointer_references(SSA *sap) {
   SSB sb(sap);
   sb.foo();
 }
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK30 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  %s --check-prefix CK30 --check-prefix CK30-64
+// RUN: %clang_cc1 -DCK30 -fopenmp -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-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefix CK30 --check-prefix CK30-64
+// RUN: %clang_cc1 -DCK30 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefix CK30 --check-prefix CK30-32
+// RUN: %clang_cc1 -DCK30 -fopenmp -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-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  %s  --check-prefix CK30 --check-prefix CK30-32
+
+// RUN: %clang_cc1 -DCK30 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY30 %s
+// RUN: %clang_cc1 -DCK30 -fopenmp-simd -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-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY30 %s
+// RUN: %clang_cc1 -DCK30 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY30 %s
+// RUN: %clang_cc1 -DCK30 -fopenmp-simd -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-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-ONLY30 %s
+// SIMD-ONLY30-NOT: {{__kmpc|__tgt}}
+#ifdef CK30
+
+// CK30-DAG: [[BASE:%.+]] = type { i32*, i32, i32* }
+// CK30-DAG: [[STRUCT:%.+]] = type { [[BASE]], i32*, i32*, i32, i32* }
+
+// CK30-LABEL: @.__omp_offloading_{{.*}}map_with_deep_copy{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
+// The first element: 0x20 - OMP_MAP_TARGET_PARAM
+// 2-4: 0x1000000000003 - OMP_MAP_MEMBER_OF(0) | OMP_MAP_TO | OMP_MAP_FROM - copies all the data in structs excluding deep-copied elements (from &s to &s.ptrBase1, from &s.ptr to &s.ptr1, from &s.ptr1 to end of s).
+// 5-6: 0x1000000000013 - OMP_MAP_MEMBER_OF(0) | OMP_MAP_PTR_AND_OBJ | OMP_MAP_TO | OMP_MAP_FROM - deep copy of the pointers + pointee.
+// CK30: [[MTYPE00:@.+]] = private {{.*}}constant [6 x i64] [i64 32, i64 281474976710659, i64 281474976710659, i64 281474976710659, i64 281474976710675, i64 281474976710675]
+
+typedef struct {
+  int *ptrBase;
+  int valBase;
+  int *ptrBase1;
+} Base;
+
+typedef struct : public Base {
+  int *ptr;
+  int *ptr2;
+  int val;
+  int *ptr1;
+} StructWithPtr;
+
+// CK30-DAG: call i32 @__tgt_target(i64 -1, i8* @.__omp_offloading_{{.*}}map_with_deep_copy{{.*}}_l{{[0-9]+}}.region_id, i32 6, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{64|32}}* [[GEPS:%.+]], i64* getelementptr inbounds ([6 x i64], [6 x i64]* [[MTYPE00]], i32 0, i32 0))
+// CK30-DAG: [[GEPS]] = getelementptr inbounds [6 x i{{64|32}}], [6 x i{{64|32}}]* [[SIZES:%.+]], i32 0, i32 0
+// CK30-DAG: [[GEPP]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS:%.+]], i32 0, i32 0
+// CK30-DAG: [[GEPBP]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BASES:%.+]], i32 0, i32 0
+
+// CK30-DAG: [[BASE_PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BASES]], i32 0, i32 0
+// CK30-DAG: [[BC:%.+]] = bitcast i8** [[BASE_PTR]] to [[STRUCT]]**
+// CK30-DAG: store [[STRUCT]]* [[S:%.+]], [[STRUCT]]** [[BC]],
+// CK30-DAG: [[PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS]], i32 0, i32 0
+// CK30-DAG: [[BC:%.+]] = bitcast i8** [[PTR]] to [[STRUCT]]**
+// CK30-DAG: store [[STRUCT]]* [[S]], [[STRUCT]]** [[BC]],
+// CK30-DAG: [[SIZE:%.+]] = getelementptr inbounds [6 x i{{64|32}}], [6 x i{{64|32}}]* [[SIZES]], i32 0, i32 0
+// CK30-64-DAG: store i64 [[S_ALLOC_SIZE:%.+]], i64* [[SIZE]],
+// CK30-32-DAG: store i32 [[S_ALLOC_SIZE32:%.+]], i32* [[SIZE]],
+// CK30-32-DAG: [[S_ALLOC_SIZE32]] = trunc i64 [[S_ALLOC_SIZE:%.+]] to i32
+// CK30-DAG: [[S_ALLOC_SIZE]] = sdiv exact i64 [[DIFF:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+// CK30-DAG: [[DIFF]] = sub i64 [[S_END_BC:%.+]], [[S_BEGIN_BC:%.+]]
+// CK30-DAG: [[S_BEGIN_BC]] = ptrtoint i8* [[S_BEGIN:%.+]] to i64
+// CK30-DAG: [[S_END_BC]] = ptrtoint i8* [[S_END:%.+]] to i64
+// CK30-DAG: [[S_BEGIN]] = bitcast [[STRUCT]]* [[S]] to i8*
+// CK30-DAG: [[S_END]] = getelementptr i8, i8* [[S_LAST:%.+]], i32 1
+// CK30-DAG: [[S_LAST]] = getelementptr i8, i8* [[S_BC:%.+]], i{{64|32}} {{55|27}}
+// CK30-DAG: [[S_BC]] = bitcast [[STRUCT]]* [[S]] to i8*
+
+// CK30-DAG: [[BASE_PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BASES]], i32 0, i32 1
+// CK30-DAG: [[BC:%.+]] = bitcast i8** [[BASE_PTR]] to [[STRUCT]]**
+// CK30-DAG: store [[STRUCT]]* [[S]], [[STRUCT]]** [[BC]],
+// CK30-DAG: [[PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS]], i32 0, i32 1
+// CK30-DAG: [[BC:%.+]] = bitcast i8** [[PTR]] to [[STRUCT]]**
+// CK30-DAG: store [[STRUCT]]* [[S]], [[STRUCT]]** [[BC]],
+// CK30-DAG: [[SIZE:%.+]] = getelementptr inbounds [6 x i{{64|32}}], [6 x i{{64|32}}]* [[SIZES]], i32 0, i32 1
+// CK30-64-DAG: store i64 [[SIZE1:%.+]], i64* [[SIZE]],
+// CK30-32-DAG: store i32 [[SIZE1_32:%.+]], i32* [[SIZE]],
+// CK30-32-DAG: [[SIZE1_32]] = trunc i64 [[SIZE1:%.+]] to i32
+// CK30-DAG: [[SIZE1]] = sdiv exact i64 [[DIFF:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+// CK30-DAG: [[DIFF]] = sub i64 [[S_PTRBASE1_BC:%.+]], [[S_BEGIN_BC:%.+]]
+// CK30-DAG: [[S_BEGIN_BC]] = ptrtoint i8* [[S_BEGIN:%.+]] to i64
+// CK30-DAG: [[S_PTRBASE1_BC]] = ptrtoint i8* [[S_PTRBASE1:%.+]] to i64
+// CK30-DAG: [[S_PTRBASE1]] = bitcast i32** [[S_PTRBASE1_REF:%.+]] to i8*
+// CK30-DAG: [[S_BEGIN]] = bitcast [[STRUCT]]* [[S]] to i8*
+// CK30-DAG: [[S_PTRBASE1_REF]] = getelementptr inbounds [[BASE]], [[BASE]]* [[BASE_ADDR:%.+]], i32 0, i32 2
+// CK30-DAG: [[BASE_ADDR]] = bitcast [[STRUCT]]* [[S]] to [[BASE]]*
+
+// CK30-DAG: [[BASE_PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BASES]], i32 0, i32 2
+// CK30-DAG: [[BC:%.+]] = bitcast i8** [[BASE_PTR]] to [[STRUCT]]**
+// CK30-DAG: store [[STRUCT]]* [[S]], [[STRUCT]]** [[BC]],
+// CK30-DAG: [[PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS]], i32 0, i32 2
+// CK30-DAG: [[BC:%.+]] = bitcast i8** [[PTR]] to i32***
+// CK30-DAG: store i32** [[PTR1:%.+]], i32*** [[BC]],
+// CK30-DAG: [[SIZE:%.+]] = getelementptr inbounds [6 x i{{64|32}}], [6 x i{{64|32}}]* [[SIZES]], i32 0, i32 2
+// CK30-64-DAG: store i64 [[SIZE2:%.+]], i64* [[SIZE]],
+// CK30-32-DAG: store i32 [[SIZE2_32:%.+]], i32* [[SIZE]],
+// CK30-32-DAG: [[SIZE2_32]] = trunc i64 [[SIZE2:%.+]] to i32
+// CK30-DAG: [[PTR1]] = getelementptr i32*, i32** [[S_PTRBASE1_REF]], i{{64|32}} 1
+// CK30-DAG: [[SIZE2]] = sdiv exact i64 [[DIFF:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+// CK30-DAG: [[DIFF]] = sub i64 [[S_PTR1_BC:%.+]], [[S_PTRBASE1_BC:%.+]]
+// CK30-DAG: [[S_PTR1_BC]] = ptrtoint i8* [[S_PTR1:%.+]] to i64
+// CK30-DAG: [[S_PTRBASE1_BC]] = ptrtoint i8* [[S_PTRBASE1:%.+]] to i64
+// CK30-DAG: [[S_PTR1]] = bitcast i32** [[S_PTR1_REF:%.+]] to i8*
+// CK30-DAG: [[S_PTRBASE1]] = bitcast i32** [[PTR1]] to i8*
+// CK30-DAG: [[S_PTR1_REF]] = getelementptr inbounds [[STRUCT]], [[STRUCT]]* [[S]], i32 0, i32 4
+
+// CK30-DAG: [[BASE_PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BASES]], i32 0, i32 3
+// CK30-DAG: [[BC:%.+]] = bitcast i8** [[BASE_PTR]] to [[STRUCT]]**
+// CK30-DAG: store [[STRUCT]]* [[S]], [[STRUCT]]** [[BC]],
+// CK30-DAG: [[PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS]], i32 0, i32 3
+// CK30-DAG: [[BC:%.+]] = bitcast i8** [[PTR]] to i32***
+// CK30-DAG: store i32** [[PTR2:%.+]], i32*** [[BC]],
+// CK30-DAG: [[SIZE:%.+]] = getelementptr inbounds [6 x i{{64|32}}], [6 x i{{64|32}}]* [[SIZES]], i32 0, i32 3
+// CK30-64-DAG: store i64 [[SIZE3:%.+]], i64* [[SIZE]],
+// CK30-32-DAG: store i32 [[SIZE3_32:%.+]], i32* [[SIZE]],
+// CK30-32-DAG: [[SIZE3_32]] = trunc i64 [[SIZE3:%.+]] to i32
+// CK30-DAG: [[PTR2]] = getelementptr i32*, i32** [[S_PTR1_REF]], i{{64|32}} 1
+// CK30-DAG: [[SIZE3]] = sdiv exact i64 [[DIFF:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+// CK30-DAG: [[DIFF]] = sub i64 [[S_END_BC:%.+]], [[S_PTR1_BC:%.+]]
+// CK30-DAG: [[S_PTR1_BC]] = ptrtoint i8* [[S_PTR1:%.+]] to i64
+// CK30-DAG: [[S_END_BC]] = ptrtoint i8* [[S_END:%.+]] to i64
+// CK30-DAG: [[S_PTR1]] = bitcast i32** [[PTR2]] to i8*
+// CK30-DAG: [[S_END]] = getelementptr i8, i8* [[S_LAST]], i{{64|32}} 1
+
+// CK30-DAG: [[BASE_PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BASES]], i32 0, i32 4
+// CK30-DAG: [[BC:%.+]] = bitcast i8** [[BASE_PTR]] to i32***
+// CK30-DAG: store i32** [[S_PTR1:%.+]], i32*** [[BC]],
+// CK30-DAG: [[PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS]], i32 0, i32 4
+// CK30-DAG: [[BC:%.+]] = bitcast i8** [[PTR]] to i32**
+// CK30-DAG: store i32* [[S_PTR1_BEGIN:%.+]], i32** [[BC]],
+// CK30-DAG: [[SIZE:%.+]] = getelementptr inbounds [6 x i{{64|32}}], [6 x i{{64|32}}]* [[SIZES]], i32 0, i32 4
+// CK30-DAG: store i{{64|32}} 4, i{{64|32}}* [[SIZE]],
+// CK30-DAG: [[S_PTR1]] = getelementptr inbounds [[STRUCT]], [[STRUCT]]* [[S]], i32 0, i32 4
+// CK30-DAG: [[S_PTR1_BEGIN]] = getelementptr inbounds i32, i32* [[S_PTR1_BEGIN_REF:%.+]], i{{64|32}} 0
+// CK30-DAG: [[S_PTR1_BEGIN_REF]] = load i32*, i32** [[S_PTR1:%.+]],
+// CK30-DAG: [[S_PTR1]] = getelementptr inbounds [[STRUCT]], [[STRUCT]]* [[S]], i32 0, i32 4
+
+// CK30-DAG: [[BASE_PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BASES]], i32 0, i32 5
+// CK30-DAG: [[BC:%.+]] = bitcast i8** [[BASE_PTR]] to i32***
+// CK30-DAG: store i32** [[S_PTRBASE1:%.+]], i32*** [[BC]],
+// CK30-DAG: [[PTR:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS]], i32 0, i32 5
+// CK30-DAG: [[BC:%.+]] = bitcast i8** [[PTR]] to i32**
+// CK30-DAG: store i32* [[S_PTRBASE1_BEGIN:%.+]], i32** [[BC]],
+// CK30-DAG: [[SIZE:%.+]] = getelementptr inbounds [6 x i{{64|32}}], [6 x i{{64|32}}]* [[SIZES]], i32 0, i32 5
+// CK30-DAG: store i{{64|32}} 4, i{{64|32}}* [[SIZE]],
+// CK30-DAG: [[S_PTRBASE1]] = getelementptr inbounds [[BASE]], [[BASE]]* [[S_BASE:%.+]], i32 0, i32 2
+// CK30-DAG: [[S_BASE]] = bitcast [[STRUCT]]* [[S]] to [[BASE]]*
+// CK30-DAG: [[S_PTRBASE1_BEGIN]] = getelementptr inbounds i32, i32* [[S_PTRBASE1_BEGIN_REF:%.+]], i{{64|32}} 0
+// CK30-DAG: [[S_PTRBASE1_BEGIN_REF]] = load i32*, i32** [[S_PTRBASE1:%.+]],
+// CK30-DAG: [[S_PTRBASE1]] = getelementptr inbounds [[BASE]], [[BASE]]* [[S_BASE:%.+]], i32 0, i32 2
+// CK30-DAG: [[S_BASE]] = bitcast [[STRUCT]]* [[S]] to [[BASE]]*
+void map_with_deep_copy() {
+  StructWithPtr s;
+#pragma omp target map(s, s.ptr1 [0:1], s.ptrBase1 [0:1])
+  {
+    s.val++;
+    s.ptr1[0]++;
+    s.ptrBase1[0] = 10001;
+  }
+}
+
 #endif
 #endif

clang/test/OpenMP/target_map_messages.cpp

@@ -591,9 +591,7 @@ int main(int argc, char **argv) {
 #pragma omp target map(s.b[:5])
 // expected-error@+1 {{variable already marked as mapped in current construct}}
   { s.a++; }
-// expected-note@+1 {{used here}}
 #pragma omp target map(s.p[:5])
-// expected-error@+1 {{variable already marked as mapped in current construct}}
   { s.a++; }
 // expected-note@+1 {{used here}}
 #pragma omp target map(s.s.sa[3].a)