Index: clang/include/clang/AST/OpenMPClause.h
===================================================================
--- clang/include/clang/AST/OpenMPClause.h
+++ clang/include/clang/AST/OpenMPClause.h
@@ -5354,7 +5354,7 @@
   /// Map-type-modifiers for the 'map' clause.
   OpenMPMapModifierKind MapTypeModifiers[NumberOfOMPMapClauseModifiers] = {
       OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown,
-      OMPC_MAP_MODIFIER_unknown};
+      OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown};
 
   /// Location of map-type-modifiers for the 'map' clause.
   SourceLocation MapTypeModifiersLoc[NumberOfOMPMapClauseModifiers];
Index: clang/include/clang/Basic/DiagnosticParseKinds.td
===================================================================
--- clang/include/clang/Basic/DiagnosticParseKinds.td
+++ clang/include/clang/Basic/DiagnosticParseKinds.td
@@ -1248,7 +1248,7 @@
 def err_omp_unknown_map_type : Error<
   "incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'">;
 def err_omp_unknown_map_type_modifier : Error<
-  "incorrect map type modifier, expected 'always', 'close', or 'mapper'">;
+  "incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'">;
 def err_omp_map_type_missing : Error<
   "missing map type">;
 def err_omp_map_type_modifier_missing : Error<
Index: clang/include/clang/Basic/OpenMPKinds.def
===================================================================
--- clang/include/clang/Basic/OpenMPKinds.def
+++ clang/include/clang/Basic/OpenMPKinds.def
@@ -124,6 +124,7 @@
 OPENMP_MAP_MODIFIER_KIND(always)
 OPENMP_MAP_MODIFIER_KIND(close)
 OPENMP_MAP_MODIFIER_KIND(mapper)
+OPENMP_MAP_MODIFIER_KIND(present)
 
 // Modifiers for 'to' clause.
 OPENMP_TO_MODIFIER_KIND(mapper)
Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7040,6 +7040,8 @@
     /// Close is a hint to the runtime to allocate memory close to
     /// the target device.
     OMP_MAP_CLOSE = 0x400,
+    /// Produce a runtime error if the data is not already allocated.
+    OMP_MAP_PRESENT = 0x800,
     /// The 16 MSBs of the flags indicate whether the entry is member of some
     /// struct/class.
     OMP_MAP_MEMBER_OF = 0xffff000000000000,
@@ -7260,6 +7262,9 @@
     if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_close)
         != MapModifiers.end())
       Bits |= OMP_MAP_CLOSE;
+    if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_present)
+        != MapModifiers.end())
+      Bits |= OMP_MAP_PRESENT;
     return Bits;
   }
 
@@ -7926,6 +7931,13 @@
     Sizes.push_back(Size);
     // Map type is always TARGET_PARAM
     Types.push_back(OMP_MAP_TARGET_PARAM);
+    // If any element has the present modifier, then make sure the runtime
+    // doesn't attempt to allocate the struct.
+    if (CurTypes.end() !=
+        llvm::find_if(CurTypes, [](OpenMPOffloadMappingFlags Type) {
+          return Type & OMP_MAP_PRESENT;
+        }))
+      Types.back() |= OMP_MAP_PRESENT;
     // Remove TARGET_PARAM flag from the first element
     (*CurTypes.begin()) &= ~OMP_MAP_TARGET_PARAM;
 
@@ -7944,7 +7956,10 @@
   /// index where it occurs is appended to the device pointers info array.
   void generateAllInfo(MapBaseValuesArrayTy &BasePointers,
                        MapValuesArrayTy &Pointers, MapValuesArrayTy &Sizes,
-                       MapFlagsArrayTy &Types) const {
+                       MapFlagsArrayTy &Types,
+                       const llvm::DenseSet<const ValueDecl *> &SkipVarSet =
+                           llvm::DenseSet<const ValueDecl *>(),
+                       bool PresentModifierOnly = false) const {
     // We have to process the component lists that relate with the same
     // declaration in a single chunk so that we can generate the map flags
     // correctly. Therefore, we organize all lists in a map.
@@ -7953,14 +7968,23 @@
     // Helper function to fill the information map for the different supported
     // clauses.
     auto &&InfoGen =
-        [&Info](const ValueDecl *D,
-                OMPClauseMappableExprCommon::MappableExprComponentListRef L,
-                OpenMPMapClauseKind MapType,
-                ArrayRef<OpenMPMapModifierKind> MapModifiers,
-                bool ReturnDevicePointer, bool IsImplicit,
-                bool ForDeviceAddr = false) {
+        [&Info, &SkipVarSet, PresentModifierOnly](
+            const ValueDecl *D,
+            OMPClauseMappableExprCommon::MappableExprComponentListRef L,
+            OpenMPMapClauseKind MapType,
+            ArrayRef<OpenMPMapModifierKind> MapModifiers,
+            bool ReturnDevicePointer, bool IsImplicit,
+            bool ForDeviceAddr = false) {
           const ValueDecl *VD =
               D ? cast<ValueDecl>(D->getCanonicalDecl()) : nullptr;
+          if (SkipVarSet.count(VD))
+            return;
+          // Skip if we're only handling map types with the "present" modifier
+          // but we don't have it here.
+          if (PresentModifierOnly &&
+              MapModifiers.end() ==
+                  llvm::find(MapModifiers, OMPC_MAP_MODIFIER_present))
+            return;
           Info[VD].emplace_back(L, MapType, MapModifiers, ReturnDevicePointer,
                                 IsImplicit, ForDeviceAddr);
         };
@@ -9465,6 +9489,7 @@
     // Get mappable expression information.
     MappableExprsHandler MEHandler(D, CGF);
     llvm::DenseMap<llvm::Value *, llvm::Value *> LambdaPointers;
+    llvm::DenseSet<const ValueDecl *> CapturedVarSet;
 
     auto RI = CS.getCapturedRecordDecl()->field_begin();
     auto CV = CapturedVars.begin();
@@ -9493,6 +9518,10 @@
         // just do a default mapping.
         MEHandler.generateInfoForCapture(CI, *CV, CurBasePointers, CurPointers,
                                          CurSizes, CurMapTypes, PartialStruct);
+        if (!CI->capturesThis())
+          CapturedVarSet.insert(CI->getCapturedVar()->getCanonicalDecl());
+        else
+          CapturedVarSet.insert(nullptr);
         if (CurBasePointers.empty())
           MEHandler.generateDefaultMapInfo(*CI, **RI, *CV, CurBasePointers,
                                            CurPointers, CurSizes, CurMapTypes);
@@ -9530,6 +9559,11 @@
     // but "declare target link" global variables.
     MEHandler.generateInfoForDeclareTargetLink(BasePointers, Pointers, Sizes,
                                                MapTypes);
+    // Map any list items in a map clause that have "present" map type modifiers
+    // but were not captures because they weren't referenced within the
+    // construct.  The runtime must check for their presence anyway.
+    MEHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes,
+                              CapturedVarSet, /*PresentModifierOnly=*/true);
 
     TargetDataInfo Info;
     // Fill up the arrays and create the arguments.
Index: clang/lib/Parse/ParseOpenMP.cpp
===================================================================
--- clang/lib/Parse/ParseOpenMP.cpp
+++ clang/lib/Parse/ParseOpenMP.cpp
@@ -3096,7 +3096,8 @@
   while (getCurToken().isNot(tok::colon)) {
     OpenMPMapModifierKind TypeModifier = isMapModifier(*this);
     if (TypeModifier == OMPC_MAP_MODIFIER_always ||
-        TypeModifier == OMPC_MAP_MODIFIER_close) {
+        TypeModifier == OMPC_MAP_MODIFIER_close ||
+        TypeModifier == OMPC_MAP_MODIFIER_present) {
       Data.MapTypeModifiers.push_back(TypeModifier);
       Data.MapTypeModifiersLoc.push_back(Tok.getLocation());
       ConsumeToken();
Index: clang/lib/Sema/SemaOpenMP.cpp
===================================================================
--- clang/lib/Sema/SemaOpenMP.cpp
+++ clang/lib/Sema/SemaOpenMP.cpp
@@ -17575,9 +17575,9 @@
     OpenMPMapClauseKind MapType, bool IsMapTypeImplicit, SourceLocation MapLoc,
     SourceLocation ColonLoc, ArrayRef<Expr *> VarList,
     const OMPVarListLocTy &Locs, ArrayRef<Expr *> UnresolvedMappers) {
-  OpenMPMapModifierKind Modifiers[] = {OMPC_MAP_MODIFIER_unknown,
-                                       OMPC_MAP_MODIFIER_unknown,
-                                       OMPC_MAP_MODIFIER_unknown};
+  OpenMPMapModifierKind Modifiers[] = {
+      OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown,
+      OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown};
   SourceLocation ModifiersLoc[NumberOfOMPMapClauseModifiers];
 
   // Process map-type-modifiers, flag errors for duplicate modifiers.
Index: clang/test/OpenMP/target_ast_print.cpp
===================================================================
--- clang/test/OpenMP/target_ast_print.cpp
+++ clang/test/OpenMP/target_ast_print.cpp
@@ -15,7 +15,7 @@
 
 template <typename T, int C>
 T tmain(T argc, T *argv) {
-  T i, j, a[20], always, close;
+  T i, j, a[20], always, close, present;
 #pragma omp target
   foo();
 #pragma omp target if (target:argc > 0)
@@ -44,6 +44,14 @@
   {close++;}
 #pragma omp target map(close,i)
   {close++;i++;}
+#pragma omp target map(present,alloc: i)
+  foo();
+#pragma omp target map(present from: i)
+  foo();
+#pragma omp target map(present)
+  {present++;}
+#pragma omp target map(present,i)
+  {present++;i++;}
 #pragma omp target nowait
   foo();
 #pragma omp target depend(in : argc, argv[i:argc], a[:])
@@ -93,6 +101,19 @@
 // OMP45-NEXT: close++;
 // OMP45-NEXT: i++;
 // OMP45-NEXT: }
+// OMP45-NEXT: #pragma omp target map(present,alloc: i)
+// OMP45-NEXT: foo()
+// OMP45-NEXT: #pragma omp target map(present,from: i)
+// OMP45-NEXT: foo()
+// OMP45-NEXT: #pragma omp target map(tofrom: present)
+// OMP45-NEXT: {
+// OMP45-NEXT: present++;
+// OMP45-NEXT: }
+// OMP45-NEXT: #pragma omp target map(tofrom: present,i)
+// OMP45-NEXT: {
+// OMP45-NEXT: present++;
+// OMP45-NEXT: i++;
+// OMP45-NEXT: }
 // OMP45-NEXT: #pragma omp target nowait
 // OMP45-NEXT: foo()
 // OMP45-NEXT: #pragma omp target depend(in : argc,argv[i:argc],a[:])
@@ -139,6 +160,19 @@
 // OMP45-NEXT: close++;
 // OMP45-NEXT: i++;
 // OMP45-NEXT: }
+// OMP45-NEXT: #pragma omp target map(present,alloc: i)
+// OMP45-NEXT: foo()
+// OMP45-NEXT: #pragma omp target map(present,from: i)
+// OMP45-NEXT: foo()
+// OMP45-NEXT: #pragma omp target map(tofrom: present)
+// OMP45-NEXT: {
+// OMP45-NEXT: present++;
+// OMP45-NEXT: }
+// OMP45-NEXT: #pragma omp target map(tofrom: present,i)
+// OMP45-NEXT: {
+// OMP45-NEXT: present++;
+// OMP45-NEXT: i++;
+// OMP45-NEXT: }
 // OMP45-NEXT: #pragma omp target nowait
 // OMP45-NEXT: foo()
 // OMP45-NEXT: #pragma omp target depend(in : argc,argv[i:argc],a[:])
@@ -185,6 +219,19 @@
 // OMP45-NEXT: close++;
 // OMP45-NEXT: i++;
 // OMP45-NEXT: }
+// OMP45-NEXT: #pragma omp target map(present,alloc: i)
+// OMP45-NEXT: foo()
+// OMP45-NEXT: #pragma omp target map(present,from: i)
+// OMP45-NEXT: foo()
+// OMP45-NEXT: #pragma omp target map(tofrom: present)
+// OMP45-NEXT: {
+// OMP45-NEXT: present++;
+// OMP45-NEXT: }
+// OMP45-NEXT: #pragma omp target map(tofrom: present,i)
+// OMP45-NEXT: {
+// OMP45-NEXT: present++;
+// OMP45-NEXT: i++;
+// OMP45-NEXT: }
 // OMP45-NEXT: #pragma omp target nowait
 // OMP45-NEXT: foo()
 // OMP45-NEXT: #pragma omp target depend(in : argc,argv[i:argc],a[:])
@@ -229,7 +276,7 @@
 
 // OMP45-LABEL: int main(int argc, char **argv) {
 int main (int argc, char **argv) {
-  int i, j, a[20], always, close;
+  int i, j, a[20], always, close, present;
 // OMP45-NEXT: int i, j, a[20]
 #pragma omp target
 // OMP45-NEXT: #pragma omp target
@@ -310,6 +357,31 @@
 // OMP45-NEXT: i++;
 // OMP45-NEXT: }
 
+#pragma omp target map(present,alloc: i)
+// OMP45-NEXT: #pragma omp target map(present,alloc: i)
+  foo();
+// OMP45-NEXT: foo();
+
+#pragma omp target map(present from: i)
+// OMP45-NEXT: #pragma omp target map(present,from: i)
+  foo();
+// OMP45-NEXT: foo();
+
+#pragma omp target map(present)
+// OMP45-NEXT: #pragma omp target map(tofrom: present)
+  {present++;}
+// OMP45-NEXT: {
+// OMP45-NEXT: present++;
+// OMP45-NEXT: }
+
+#pragma omp target map(present,i)
+// OMP45-NEXT: #pragma omp target map(tofrom: present,i)
+  {present++;i++;}
+// OMP45-NEXT: {
+// OMP45-NEXT: present++;
+// OMP45-NEXT: i++;
+// OMP45-NEXT: }
+
 #pragma omp target nowait
 // OMP45-NEXT: #pragma omp target nowait
   foo();
@@ -363,7 +435,7 @@
 
 template <typename T, int C>
 T tmain(T argc, T *argv) {
-  T i, j, a[20], always, close;
+  T i, j, a[20], always, close, present;
 #pragma omp target device(argc)
   foo();
 #pragma omp target if (target:argc > 0) device(device_num: C)
@@ -392,6 +464,14 @@
   {close++;}
 #pragma omp target map(close,i)
   {close++;i++;}
+#pragma omp target map(present,alloc: i)
+  foo();
+#pragma omp target map(present from: i)
+  foo();
+#pragma omp target map(present)
+  {present++;}
+#pragma omp target map(present,i)
+  {present++;i++;}
 #pragma omp target nowait
   foo();
 #pragma omp target depend(in : argc, argv[i:argc], a[:])
@@ -507,6 +587,19 @@
 // OMP5-NEXT: close++;
 // OMP5-NEXT: i++;
 // OMP5-NEXT: }
+// OMP5-NEXT: #pragma omp target map(present,alloc: i)
+// OMP5-NEXT: foo()
+// OMP5-NEXT: #pragma omp target map(present,from: i)
+// OMP5-NEXT: foo()
+// OMP5-NEXT: #pragma omp target map(tofrom: present)
+// OMP5-NEXT: {
+// OMP5-NEXT: present++;
+// OMP5-NEXT: }
+// OMP5-NEXT: #pragma omp target map(tofrom: present,i)
+// OMP5-NEXT: {
+// OMP5-NEXT: present++;
+// OMP5-NEXT: i++;
+// OMP5-NEXT: }
 // OMP5-NEXT: #pragma omp target nowait
 // OMP5-NEXT: foo()
 // OMP5-NEXT: #pragma omp target depend(in : argc,argv[i:argc],a[:])
@@ -603,6 +696,19 @@
 // OMP5-NEXT: close++;
 // OMP5-NEXT: i++;
 // OMP5-NEXT: }
+// OMP5-NEXT: #pragma omp target map(present,alloc: i)
+// OMP5-NEXT: foo()
+// OMP5-NEXT: #pragma omp target map(present,from: i)
+// OMP5-NEXT: foo()
+// OMP5-NEXT: #pragma omp target map(tofrom: present)
+// OMP5-NEXT: {
+// OMP5-NEXT: present++;
+// OMP5-NEXT: }
+// OMP5-NEXT: #pragma omp target map(tofrom: present,i)
+// OMP5-NEXT: {
+// OMP5-NEXT: present++;
+// OMP5-NEXT: i++;
+// OMP5-NEXT: }
 // OMP5-NEXT: #pragma omp target nowait
 // OMP5-NEXT: foo()
 // OMP5-NEXT: #pragma omp target depend(in : argc,argv[i:argc],a[:])
@@ -699,6 +805,19 @@
 // OMP5-NEXT: close++;
 // OMP5-NEXT: i++;
 // OMP5-NEXT: }
+// OMP5-NEXT: #pragma omp target map(present,alloc: i)
+// OMP5-NEXT: foo()
+// OMP5-NEXT: #pragma omp target map(present,from: i)
+// OMP5-NEXT: foo()
+// OMP5-NEXT: #pragma omp target map(tofrom: present)
+// OMP5-NEXT: {
+// OMP5-NEXT: present++;
+// OMP5-NEXT: }
+// OMP5-NEXT: #pragma omp target map(tofrom: present,i)
+// OMP5-NEXT: {
+// OMP5-NEXT: present++;
+// OMP5-NEXT: i++;
+// OMP5-NEXT: }
 // OMP5-NEXT: #pragma omp target nowait
 // OMP5-NEXT: foo()
 // OMP5-NEXT: #pragma omp target depend(in : argc,argv[i:argc],a[:])
@@ -806,7 +925,7 @@
 
 // OMP5-LABEL: int main(int argc, char **argv) {
 int main (int argc, char **argv) {
-  int i, j, a[20], always, close;
+  int i, j, a[20], always, close, present;
 // OMP5-NEXT: int i, j, a[20]
 #pragma omp target
 // OMP5-NEXT: #pragma omp target
@@ -887,6 +1006,31 @@
 // OMP5-NEXT: i++;
 // OMP5-NEXT: }
 
+#pragma omp target map(present,alloc: i)
+// OMP5-NEXT: #pragma omp target map(present,alloc: i)
+  foo();
+// OMP5-NEXT: foo();
+
+#pragma omp target map(present from: i)
+// OMP5-NEXT: #pragma omp target map(present,from: i)
+  foo();
+// OMP5-NEXT: foo();
+
+#pragma omp target map(present)
+// OMP5-NEXT: #pragma omp target map(tofrom: present)
+  {present++;}
+// OMP5-NEXT: {
+// OMP5-NEXT: present++;
+// OMP5-NEXT: }
+
+#pragma omp target map(present,i)
+// OMP5-NEXT: #pragma omp target map(tofrom: present,i)
+  {present++;i++;}
+// OMP5-NEXT: {
+// OMP5-NEXT: present++;
+// OMP5-NEXT: i++;
+// OMP5-NEXT: }
+
 #pragma omp target nowait
 // OMP5-NEXT: #pragma omp target nowait
   foo();
Index: clang/test/OpenMP/target_data_ast_print.cpp
===================================================================
--- clang/test/OpenMP/target_data_ast_print.cpp
+++ clang/test/OpenMP/target_data_ast_print.cpp
@@ -43,6 +43,9 @@
 #pragma omp target data map(close,alloc: e)
   foo();
 
+#pragma omp target data map(present,alloc: e)
+  foo();
+
 // nesting a target region
 #pragma omp target data map(e)
 {
@@ -50,6 +53,8 @@
     foo();
   #pragma omp target map(close, alloc: e)
     foo();
+  #pragma omp target map(present, alloc: e)
+    foo();
 }
 
   return 0;
@@ -75,12 +80,16 @@
 // CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target data map(close,alloc: e)
 // CHECK-NEXT: foo();
+// CHECK-NEXT: #pragma omp target data map(present,alloc: e)
+// CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target data map(tofrom: e)
 // CHECK-NEXT: {
 // CHECK-NEXT: #pragma omp target map(always,alloc: e)
 // CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target map(close,alloc: e)
 // CHECK-NEXT: foo();
+// CHECK-NEXT: #pragma omp target map(present,alloc: e)
+// CHECK-NEXT: foo();
 // CHECK: template<> int tmain<int, 5>(int argc, int *argv) {
 // CHECK-NEXT: int i, j, b, c, d, e, x[20];
 // CHECK-NEXT: #pragma omp target data map(to: c)
@@ -101,12 +110,16 @@
 // CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target data map(close,alloc: e)
 // CHECK-NEXT: foo();
+// CHECK-NEXT: #pragma omp target data map(present,alloc: e)
+// CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target data map(tofrom: e)
 // CHECK-NEXT: {
 // CHECK-NEXT: #pragma omp target map(always,alloc: e)
 // CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target map(close,alloc: e)
 // CHECK-NEXT: foo();
+// CHECK-NEXT: #pragma omp target map(present,alloc: e)
+// CHECK-NEXT: foo();
 // CHECK: template<> char tmain<char, 1>(char argc, char *argv) {
 // CHECK-NEXT: char i, j, b, c, d, e, x[20];
 // CHECK-NEXT: #pragma omp target data map(to: c)
@@ -127,12 +140,16 @@
 // CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target data map(close,alloc: e)
 // CHECK-NEXT: foo();
+// CHECK-NEXT: #pragma omp target data map(present,alloc: e)
+// CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target data map(tofrom: e)
 // CHECK-NEXT: {
 // CHECK-NEXT: #pragma omp target map(always,alloc: e)
 // CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target map(close,alloc: e)
 // CHECK-NEXT: foo();
+// CHECK-NEXT: #pragma omp target map(present,alloc: e)
+// CHECK-NEXT: foo();
 
 int main (int argc, char **argv) {
   int b = argc, c, d, e, f, g, x[20];
@@ -185,6 +202,11 @@
   foo();
 // CHECK-NEXT: foo();
 
+#pragma omp target data map(present,alloc: e)
+// CHECK-NEXT: #pragma omp target data map(present,alloc: e)
+  foo();
+// CHECK-NEXT: foo();
+
 // nesting a target region
 #pragma omp target data map(e)
 // CHECK-NEXT: #pragma omp target data map(tofrom: e)
@@ -197,6 +219,10 @@
 #pragma omp target map(close, alloc: e)
 // CHECK-NEXT: #pragma omp target map(close,alloc: e)
   foo();
+// CHECK-NEXT: foo();
+#pragma omp target map(always, alloc: e)
+// CHECK-NEXT: #pragma omp target map(always,alloc: e)
+  foo();
 }
 
   return tmain<int, 5>(argc, &argc) + tmain<char, 1>(argv[0][0], argv[0]);
Index: clang/test/OpenMP/target_data_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_data_codegen.cpp
+++ clang/test/OpenMP/target_data_codegen.cpp
@@ -44,6 +44,10 @@
 
 // CK1: [[MTYPE06:@.+]] = {{.+}}constant [1 x i64] [i64 1061]
 
+// CK1: [[MTYPE07:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x821]]]
+
+// CK1: [[MTYPE08:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0xC25]]]
+
 // CK1-LABEL: _Z3fooi
 void foo(int arg) {
   int la;
@@ -225,6 +229,59 @@
   // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S]]
   #pragma omp target data map(always close, to: lb)
   {++arg;}
+
+  // Region 07
+  // CK1-DAG: call void @__tgt_target_data_begin(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE07]]{{.+}})
+  // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+  // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+
+  // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+  // CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
+  // CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
+  // CK1-DAG: store float* [[VAR0:%.+]], float** [[CBP0]]
+  // CK1-DAG: store float* [[VAR0]], float** [[CP0]]
+  // CK1-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
+  // CK1-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
+  // CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
+  // CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
+  // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
+
+  // CK1-DAG: call void @__tgt_target_data_end(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE07]]{{.+}})
+  // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
+  // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
+  // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S]]
+  #pragma omp target data map(present, to: lb)
+  {++arg;}
+
+  // Region 08
+  // CK1-DAG: call void @__tgt_target_data_begin(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE08]]{{.+}})
+  // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+  // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+
+  // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+  // CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
+  // CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
+  // CK1-DAG: store float* [[VAR0:%.+]], float** [[CBP0]]
+  // CK1-DAG: store float* [[VAR0]], float** [[CP0]]
+  // CK1-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
+  // CK1-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
+  // CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
+  // CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
+  // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
+
+  // CK1-DAG: call void @__tgt_target_data_end(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE08]]{{.+}})
+  // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
+  // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
+  // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S]]
+  #pragma omp target data map(always close present, to: lb)
+  {++arg;}
+
 }
 #endif
 ///==========================================================================///
@@ -370,13 +427,13 @@
 
   T foo(T arg) {
     // Region 00
-    #pragma omp target data map(always, close to: b[1:3]) if(a>123) device(arg)
+    #pragma omp target data map(always, close, present to: b[1:3]) if(a>123) device(arg)
     {arg++;}
     return arg;
   }
 };
 
-// CK4: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976711701]
+// CK4: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 [[#0x820]], i64 [[#0x1000000000C15]]]
 
 // CK4-LABEL: _Z3bari
 int bar(int arg){
@@ -491,4 +548,77 @@
   {++arg;}
 }
 #endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK7 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK7 --check-prefix CK7-64
+// RUN: %clang_cc1 -DCK7 -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 %s  --check-prefix CK7 --check-prefix CK7-64
+// RUN: %clang_cc1 -DCK7 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK7 --check-prefix CK7-32
+// RUN: %clang_cc1 -DCK7 -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 %s  --check-prefix CK7 --check-prefix CK7-32
+
+// RUN: %clang_cc1 -DCK7 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
+// RUN: %clang_cc1 -DCK7 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
+// RUN: %clang_cc1 -DCK7 -verify -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
+// RUN: %clang_cc1 -DCK7 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
+// SIMD-ONLY2-NOT: {{__kmpc|__tgt}}#ifdef CK7
+#ifdef CK7
+struct S1 {
+  int i;
+};
+struct S2 {
+  S1 s;
+  struct S2 *ps;
+};
+
+void test_present_modifier(int arg) {
+  S2 *ps1;
+  S2 *ps2;
+
+  // Make sure the struct picks up present even if another element of the struct
+  // doesn't have present.
+
+  // CK7: private unnamed_addr constant [11 x i64]
+
+  // ps1
+  // CK7-SAME: {{^}} [i64 [[#0x820]], i64 [[#0x1000000000003]],
+  // CK7-SAME: {{^}} i64 [[#0x1000000000810]], i64 [[#0x810]], i64 [[#0x813]],
+
+  // arg
+  // CK7-SAME: {{^}} i64 [[#0x823]],
+
+  // ps2
+  // CK7-SAME: {{^}} i64 [[#0x820]], i64 [[#0x7000000000803]],
+  // CK7-SAME: {{^}} i64 [[#0x7000000000010]], i64 [[#0x10]], i64 [[#0x13]]]
+  #pragma omp target data map(tofrom: ps1->s) \
+                          map(present,tofrom: arg, ps1->ps->ps->ps->s, ps2->s) \
+                          map(tofrom: ps2->ps->ps->ps->s)
+  {
+    ++(arg);
+  }
+}
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK8 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK8 --check-prefix CK8-64
+// RUN: %clang_cc1 -DCK8 -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 %s  --check-prefix CK8 --check-prefix CK8-64
+// RUN: %clang_cc1 -DCK8 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK8 --check-prefix CK8-32
+// RUN: %clang_cc1 -DCK8 -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 %s  --check-prefix CK8 --check-prefix CK8-32
+
+// RUN: %clang_cc1 -DCK8 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
+// RUN: %clang_cc1 -DCK8 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
+// RUN: %clang_cc1 -DCK8 -verify -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
+// RUN: %clang_cc1 -DCK8 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
+// SIMD-ONLY2-NOT: {{__kmpc|__tgt}}
+#ifdef CK8
+void test_present_modifier(int arg) {
+  // CK8: private unnamed_addr constant [1 x i64] [i64 [[#0x823]]]
+  #pragma omp target data map(present,tofrom: arg)
+  {++arg;}
+}
+#endif
 #endif
Index: clang/test/OpenMP/target_defaultmap_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_defaultmap_codegen.cpp
+++ clang/test/OpenMP/target_defaultmap_codegen.cpp
@@ -1419,7 +1419,7 @@
 
 // CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK24: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
-// CK24: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 1063]
+// CK24: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 [[#0xC27]]]
 
 // CK24-LABEL: explicit_maps_single{{.*}}(
 void explicit_maps_single (int ii){
@@ -1445,7 +1445,7 @@
    a++;
   }
 
-  // Always Close.
+  // Always Close Present.
   // Region 01
   // CK24-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
   // CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
@@ -1459,7 +1459,7 @@
   // CK24-DAG: store i32* [[VAR0]], i32** [[CP0]]
 
   // CK24: call void [[CALL01:@.+]](i32* {{[^,]+}})
-  #pragma omp target map(always close tofrom: a) defaultmap(none:scalar)
+  #pragma omp target map(always close present tofrom: a) defaultmap(none:scalar)
   {
    a++;
   }
Index: clang/test/OpenMP/target_enter_data_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_enter_data_codegen.cpp
+++ clang/test/OpenMP/target_enter_data_codegen.cpp
@@ -42,7 +42,9 @@
 
 // CK1: [[MTYPE05:@.+]] = {{.+}}constant [1 x i64] [i64 1057]
 
-// CK1: [[MTYPE06:@.+]] = {{.+}}constant [1 x i64] [i64 1061]
+// CK1: [[MTYPE06:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x821]]]
+
+// CK1: [[MTYPE07:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0xC25]]]
 
 // CK1-LABEL: _Z3fooi
 void foo(int arg) {
@@ -176,11 +178,33 @@
   #pragma omp target enter data map(close, to: lb)
   {++arg;}
 
+  // Region 06
+  // CK1-DAG: call void @__tgt_target_data_begin(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE06]]{{.+}})
+  // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+  // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+
+  // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+  // CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
+  // CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
+  // CK1-DAG: store float* [[VAR0:%.+]], float** [[CBP0]]
+  // CK1-DAG: store float* [[VAR0]], float** [[CP0]]
+  // CK1-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
+  // CK1-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
+  // CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
+  // CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
+  // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
+  // CK1-NOT: __tgt_target_data_end
+  #pragma omp target enter data map(present, to: lb)
+  {++arg;}
+
   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
   {++arg;}
 
-  // Region 06
-  // CK1-DAG: call void @__tgt_target_data_begin(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE06]]{{.+}})
+  // Region 07
+  // CK1-DAG: call void @__tgt_target_data_begin(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE07]]{{.+}})
   // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
   // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
@@ -198,7 +222,7 @@
   // CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
   // CK1-NOT: __tgt_target_data_end
-  #pragma omp target enter data map(always close, to: lb)
+  #pragma omp target enter data map(always close present, to: lb)
   {++arg;}
 }
 #endif
@@ -377,13 +401,13 @@
 
   T foo(T arg) {
     // Region 00
-    #pragma omp target enter data map(always close to: b[1:3]) if(a>123) device(arg)
+    #pragma omp target enter data map(always close present to: b[1:3]) if(a>123) device(arg)
     {arg++;}
     return arg;
   }
 };
 
-// CK5: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976711701]
+// CK5: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 [[#0x820]], i64 [[#0x1000000000C15]]]
 
 // CK5-LABEL: _Z3bari
 int bar(int arg){
Index: clang/test/OpenMP/target_map_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_map_codegen.cpp
+++ clang/test/OpenMP/target_map_codegen.cpp
@@ -5299,6 +5299,12 @@
 // SIMD-ONLY18-NOT: {{__kmpc|__tgt}}
 #ifdef CK31
 
+// CK31: [[ST:%.+]] = type { i32, i32 }
+struct ST {
+  int i;
+  int j;
+};
+
 // CK31-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK31: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
 // CK31: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i64] [i64 1059]
@@ -5307,11 +5313,33 @@
 // CK31: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
 // CK31: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 1063]
 
+// CK31-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
+// CK31: [[MTYPE02:@.+]] = private {{.*}}constant [7 x i64] [i64 [[#0x820]],
+// CK31-SAME: {{^}} i64 [[#0x1000000000003]], i64 [[#0x1000000000803]], i64 [[#0x823]],
+// CK31-SAME: {{^}} i64 [[#0x820]], i64 [[#0x5000000000803]], i64 [[#0x5000000000003]]]
+
+// CK31-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
+// CK31: [[MTYPE03:@.+]] = private {{.*}}constant [5 x i64] [i64 [[#0x823]],
+// CK31-SAME: {{^}} i64 [[#0x820]], i64 [[#0x2000000000803]], i64 [[#0x820]],
+// CK31-SAME: {{^}} i64 [[#0x4000000000803]]]
+
+// CK31-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
+// CK31: [[SIZE04:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
+// CK31: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i64] [i64 [[#0xC27]]]
+
 // CK31-LABEL: explicit_maps_single{{.*}}(
 void explicit_maps_single (int ii){
+  // CK31: alloca i32
+
   // Map of a scalar.
+  // CK31: [[A:%.+]] = alloca i32
   int a = ii;
 
+  // CK31: [[ST1:%.+]] = alloca [[ST]]
+  // CK31: [[ST2:%.+]] = alloca [[ST]]
+  struct ST st1;
+  struct ST st2;
+
   // Close.
   // Region 00
   // CK31-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
@@ -5349,9 +5377,316 @@
   {
    a++;
   }
+
+  // Present and accessed in construct.  Make sure the struct picks up present
+  // even if another element of the struct doesn't have present.
+  // Region 02
+  // CK31: [[ST1_I:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST1]], i{{.+}} 0, i{{.+}} 0
+  // CK31: [[ST1_J:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST1]], i{{.+}} 0, i{{.+}} 1
+  // CK31: [[ST2_I:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST2]], i{{.+}} 0, i{{.+}} 0
+  // CK31: [[ST2_J:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST2]], i{{.+}} 0, i{{.+}} 1
+  // CK31-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 7, i8** [[GEPBP:%[0-9]+]], i8** [[GEPP:%[0-9]+]], i64* [[GEPS:%.+]], i64* getelementptr {{.+}}[7 x i{{.+}}]* [[MTYPE02]]{{.+}})
+  // CK31-DAG: [[GEPS]] = getelementptr inbounds [7 x i64], [7 x i64]* [[S:%.+]], i{{.+}} 0, i{{.+}} 0
+  // CK31-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK31-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // st1
+  // CK31-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK31-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK31-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+  // CK31-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
+  // CK31-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
+  // CK31-DAG: store [[ST]]* [[ST1]], [[ST]]** [[CBP0]]
+  // CK31-DAG: store i32* [[ST1_I]], i32** [[CP0]]
+  // CK31-DAG: store i64 %{{.+}}, i64* [[S0]]
+
+  // st1.i
+  // CK31-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+  // CK31-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK31-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
+  // CK31-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]**
+  // CK31-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32**
+  // CK31-DAG: store [[ST]]* [[ST1]], [[ST]]** [[CBP1]]
+  // CK31-DAG: store i32* [[ST1_I]], i32** [[CP1]]
+  // CK31-DAG: store i64 4, i64* [[S1]]
+
+  // st1.j
+  // CK31-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
+  // CK31-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
+  // CK31-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2
+  // CK31-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to [[ST]]**
+  // CK31-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to i32**
+  // CK31-DAG: store [[ST]]* [[ST1]], [[ST]]** [[CBP2]]
+  // CK31-DAG: store i32* [[ST1_J]], i32** [[CP2]]
+  // CK31-DAG: store i64 4, i64* [[S2]]
+
+  // a
+  // CK31-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3
+  // CK31-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3
+  // CK31-DAG: [[S3:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 3
+  // CK31-DAG: [[CBP3:%.+]] = bitcast i8** [[BP3]] to i32**
+  // CK31-DAG: [[CP3:%.+]] = bitcast i8** [[P3]] to i32**
+  // CK31-DAG: store i32* [[A]], i32** [[CBP3]]
+  // CK31-DAG: store i32* [[A]], i32** [[CP3]]
+  // CK31-DAG: store i64 4, i64* [[S3]]
+
+  // st2
+  // CK31-DAG: [[BP4:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 4
+  // CK31-DAG: [[P4:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 4
+  // CK31-DAG: [[S4:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 4
+  // CK31-DAG: [[CBP4:%.+]] = bitcast i8** [[BP4]] to [[ST]]**
+  // CK31-DAG: [[CP4:%.+]] = bitcast i8** [[P4]] to i32**
+  // CK31-DAG: store [[ST]]* [[ST2]], [[ST]]** [[CBP4]]
+  // CK31-DAG: store i32* [[ST2_I]], i32** [[CP4]]
+  // CK31-DAG: store i64 %{{.+}}, i64* [[S4]]
+
+  // st2.i
+  // CK31-DAG: [[BP5:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 5
+  // CK31-DAG: [[P5:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 5
+  // CK31-DAG: [[S5:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 5
+  // CK31-DAG: [[CBP5:%.+]] = bitcast i8** [[BP5]] to [[ST]]**
+  // CK31-DAG: [[CP5:%.+]] = bitcast i8** [[P5]] to i32**
+  // CK31-DAG: store [[ST]]* [[ST2]], [[ST]]** [[CBP5]]
+  // CK31-DAG: store i32* [[ST2_I]], i32** [[CP5]]
+  // CK31-DAG: store i64 4, i64* [[S5]]
+
+  // st2.j
+  // CK31-DAG: [[BP6:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 6
+  // CK31-DAG: [[P6:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 6
+  // CK31-DAG: [[S6:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 6
+  // CK31-DAG: [[CBP6:%.+]] = bitcast i8** [[BP6]] to [[ST]]**
+  // CK31-DAG: [[CP6:%.+]] = bitcast i8** [[P6]] to i32**
+  // CK31-DAG: store [[ST]]* [[ST2]], [[ST]]** [[CBP6]]
+  // CK31-DAG: store i32* [[ST2_J]], i32** [[CP6]]
+  // CK31-DAG: store i64 4, i64* [[S6]]
+
+  // CK31: call void [[CALL02:@.+]]([[ST]]* [[ST1]], i32* [[A]], [[ST]]* [[ST2]])
+  #pragma omp target map(tofrom: st1.i) map(present, tofrom: a, st1.j, st2.i) map(tofrom: st2.j)
+  {
+    st1.i++;
+    a++;
+    st1.j++;
+    st2.i++;
+    st2.j++;
+  }
+
+  // Present and not accessed in construct.  Make sure the struct picks up
+  // present even if another element of the struct doesn't have present.
+  // Region 03
+  // CK31: [[ST1_J:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST1]], i{{.+}} 0, i{{.+}} 1
+  // CK31: [[ST2_I:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST2]], i{{.+}} 0, i{{.+}} 0
+  // CK31-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 5, i8** [[GEPBP:%[0-9]+]], i8** [[GEPP:%[0-9]+]], i64* [[GEPS:%.+]], i64* getelementptr {{.+}}[5 x i{{.+}}]* [[MTYPE03]]{{.+}})
+  // CK31-DAG: [[GEPS]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S:%.+]], i{{.+}} 0, i{{.+}} 0
+  // CK31-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK31-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // a
+  // CK31-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK31-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK31-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+  // CK31-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**
+  // CK31-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
+  // CK31-DAG: store i32* [[A]], i32** [[CBP0]]
+  // CK31-DAG: store i32* [[A]], i32** [[CP0]]
+  // CK31-DAG: store i64 4, i64* [[S0]]
+
+  // st1
+  // CK31-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+  // CK31-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK31-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
+  // CK31-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]**
+  // CK31-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32**
+  // CK31-DAG: store [[ST]]* [[ST1]], [[ST]]** [[CBP1]]
+  // CK31-DAG: store i32* [[ST1_J]], i32** [[CP1]]
+  // CK31-DAG: store i64 %{{.+}}, i64* [[S1]]
+
+  // st1.j
+  // CK31-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
+  // CK31-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
+  // CK31-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2
+  // CK31-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to [[ST]]**
+  // CK31-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to i32**
+  // CK31-DAG: store [[ST]]* [[ST1]], [[ST]]** [[CBP2]]
+  // CK31-DAG: store i32* [[ST1_J]], i32** [[CP2]]
+  // CK31-DAG: store i64 4, i64* [[S2]]
+
+  // st2
+  // CK31-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3
+  // CK31-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3
+  // CK31-DAG: [[S3:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 3
+  // CK31-DAG: [[CBP3:%.+]] = bitcast i8** [[BP3]] to [[ST]]**
+  // CK31-DAG: [[CP3:%.+]] = bitcast i8** [[P3]] to i32**
+  // CK31-DAG: store [[ST]]* [[ST2]], [[ST]]** [[CBP3]]
+  // CK31-DAG: store i32* [[ST2_I]], i32** [[CP3]]
+  // CK31-DAG: store i64 %{{.+}}, i64* [[S3]]
+
+  // st2.i
+  // CK31-DAG: [[BP4:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 4
+  // CK31-DAG: [[P4:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 4
+  // CK31-DAG: [[S4:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 4
+  // CK31-DAG: [[CBP4:%.+]] = bitcast i8** [[BP4]] to [[ST]]**
+  // CK31-DAG: [[CP4:%.+]] = bitcast i8** [[P4]] to i32**
+  // CK31-DAG: store [[ST]]* [[ST2]], [[ST]]** [[CBP4]]
+  // CK31-DAG: store i32* [[ST2_I]], i32** [[CP4]]
+  // CK31-DAG: store i64 4, i64* [[S4]]
+
+  // CK31: call void [[CALL03:@.+]]()
+  #pragma omp target map(tofrom: st1.i) map(present, tofrom: a, st1.j, st2.i) map(tofrom: st2.j)
+  {
+  }
+
+  // Always Close Present.
+  // Region 04
+  // CK31-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE04]]{{.+}})
+  // CK31-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK31-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK31-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK31-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK31-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**
+  // CK31-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
+  // CK31-DAG: store i32* [[VAR0:%.+]], i32** [[CBP0]]
+  // CK31-DAG: store i32* [[VAR0]], i32** [[CP0]]
+
+  // CK31: call void [[CALL04:@.+]](i32* {{[^,]+}})
+  #pragma omp target map(always close present tofrom: a)
+  {
+   a++;
+  }
+
 }
 // CK31: define {{.+}}[[CALL00]]
 // CK31: define {{.+}}[[CALL01]]
+// CK31: define {{.+}}[[CALL02]]
+// CK31: define {{.+}}[[CALL03]]
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK31A -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 CK31A --check-prefix CK31A-64
+// RUN: %clang_cc1 -DCK31A -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 CK31A --check-prefix CK31A-64
+// RUN: %clang_cc1 -DCK31A -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 CK31A --check-prefix CK31A-32
+// RUN: %clang_cc1 -DCK31A -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 CK31A --check-prefix CK31A-32
+
+// RUN: %clang_cc1 -DCK31A -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-ONLY18 %s
+// RUN: %clang_cc1 -DCK31A -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-ONLY18 %s
+// RUN: %clang_cc1 -DCK31A -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-ONLY18 %s
+// RUN: %clang_cc1 -DCK31A -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-ONLY18 %s
+// SIMD-ONLY18-NOT: {{__kmpc|__tgt}}
+#ifdef CK31A
+
+// CK31A: [[ST:%.+]] = type { i32, i32 }
+
+// CK31A-LABEL: @.__omp_offloading_{{.*}}test_present_members{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
+// CK31A: [[MTYPE00:@.+]] = private {{.*}}constant [3 x i64] [i64 [[#0x820]],
+// CK31A-SAME: {{^}} i64 [[#0x1000000000003]], i64 [[#0x1000000000803]]]
+
+// CK31A-LABEL: @.__omp_offloading_{{.*}}test_present_members{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
+// CK31A: [[MTYPE01:@.+]] = private {{.*}}constant [2 x i64] [i64 [[#0x820]],
+// CK31A-SAME: {{^}} i64 [[#0x1000000000803]]]
+
+struct ST {
+  int i;
+  int j;
+  // CK31A-LABEL: define {{.*}}test_present_members00{{.*}}(
+  void test_present_members00() {
+    // Present and accessed in construct.  Make sure the struct picks up present
+    // even if another element of the struct doesn't have present.
+    // Region 00
+    // CK31A: [[I:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[THIS:%.+]], i{{.+}} 0, i{{.+}} 0
+    // CK31A: [[J:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[THIS]], i{{.+}} 0, i{{.+}} 1
+    // CK31A-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 3, i8** [[GEPBP:%[0-9]+]], i8** [[GEPP:%[0-9]+]], i64* [[GEPS:%.+]], i64* getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE00]]{{.+}})
+    // CK31A-DAG: [[GEPS]] = getelementptr inbounds [3 x i64], [3 x i64]* [[S:%.+]], i{{.+}} 0, i{{.+}} 0
+    // CK31A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+    // CK31A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+    // this
+    // CK31A-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+    // CK31A-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+    // CK31A-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+    // CK31A-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
+    // CK31A-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
+    // CK31A-DAG: store [[ST]]* [[THIS]], [[ST]]** [[CBP0]]
+    // CK31A-DAG: store i32* [[I]], i32** [[CP0]]
+    // CK31A-DAG: store i64 %{{.+}}, i64* [[S0]]
+
+    // i
+    // CK31A-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+    // CK31A-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+    // CK31A-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
+    // CK31A-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]**
+    // CK31A-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32**
+    // CK31A-DAG: store [[ST]]* [[THIS]], [[ST]]** [[CBP1]]
+    // CK31A-DAG: store i32* [[I]], i32** [[CP1]]
+    // CK31A-DAG: store i64 4, i64* [[S1]]
+
+    // j
+    // CK31A-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
+    // CK31A-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
+    // CK31A-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2
+    // CK31A-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to [[ST]]**
+    // CK31A-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to i32**
+    // CK31A-DAG: store [[ST]]* [[THIS]], [[ST]]** [[CBP2]]
+    // CK31A-DAG: store i32* [[J]], i32** [[CP2]]
+    // CK31A-DAG: store i64 4, i64* [[S2]]
+
+    // CK31A: call void [[CALL00:@.+]]([[ST]]* [[THIS]])
+    #pragma omp target map(tofrom: i) map(present, tofrom: j)
+    {
+      i++;
+      j++;
+    }
+  }
+
+  // CK31A-LABEL: define {{.*}}test_present_members01{{.*}}(
+  void test_present_members01() {
+    // Present and not accessed in construct.  Make sure the struct picks up
+    // present even if another element of the struct doesn't have present.
+    // Region 01
+    // CK31A: [[J:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[THIS]], i{{.+}} 0, i{{.+}} 1
+    // CK31A-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%[0-9]+]], i8** [[GEPP:%[0-9]+]], i64* [[GEPS:%.+]], i64* getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE01]]{{.+}})
+    // CK31A-DAG: [[GEPS]] = getelementptr inbounds [2 x i64], [2 x i64]* [[S:%.+]], i{{.+}} 0, i{{.+}} 0
+    // CK31A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+    // CK31A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+    // this
+    // CK31A-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+    // CK31A-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+    // CK31A-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+    // CK31A-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
+    // CK31A-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
+    // CK31A-DAG: store [[ST]]* [[THIS]], [[ST]]** [[CBP0]]
+    // CK31A-DAG: store i32* [[J]], i32** [[CP0]]
+    // CK31A-DAG: store i64 %{{.+}}, i64* [[S0]]
+
+    // j
+    // CK31A-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+    // CK31A-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+    // CK31A-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
+    // CK31A-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]**
+    // CK31A-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32**
+    // CK31A-DAG: store [[ST]]* [[THIS]], [[ST]]** [[CBP1]]
+    // CK31A-DAG: store i32* [[J]], i32** [[CP1]]
+    // CK31A-DAG: store i64 4, i64* [[S1]]
+
+    // CK31A: call void [[CALL01:@.+]]()
+    #pragma omp target map(tofrom: i) map(present, tofrom: j)
+    {
+    }
+  }
+};
+
+void test() {
+  ST s;
+  s.test_present_members00();
+  s.test_present_members01();
+}
+
+// CK31A: define {{.+}}[[CALL00]]
+// CK31A: define {{.+}}[[CALL01]]
 
 #endif
 ///==========================================================================///
Index: clang/test/OpenMP/target_map_messages.cpp
===================================================================
--- clang/test/OpenMP/target_map_messages.cpp
+++ clang/test/OpenMP/target_map_messages.cpp
@@ -116,9 +116,21 @@
     {}
     #pragma omp target map(close)   // expected-error {{use of undeclared identifier 'close'}}
     {}
+    #pragma omp target map(present, tofrom: c,f)
+    {}
+    #pragma omp target map(present, tofrom: c[1:2],f)
+    {}
+    #pragma omp target map(present, tofrom: c,f[1:2])
+    {}
+    #pragma omp target map(present, tofrom: c[:],f)   // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+    {}
+    #pragma omp target map(present, tofrom: c,f[:])   // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+    {}
+    #pragma omp target map(present)   // expected-error {{use of undeclared identifier 'present'}}
+    {}
     #pragma omp target map(close, close, tofrom: a)   // expected-error {{same map type modifier has been specified more than once}}
     {}
-    #pragma omp target map(always, close, always, close, tofrom: a)   // expected-error {{same map type modifier has been specified more than once}} expected-error {{same map type modifier has been specified more than once}}
+    #pragma omp target map(always, close, present, always, close, present, tofrom: a)   // expected-error 3 {{same map type modifier has been specified more than once}}
     {}
     #pragma omp target map( , tofrom: a)   // expected-error {{missing map type modifier}}
     {}
@@ -126,17 +138,17 @@
     {}
     #pragma omp target map( , , : a)   // expected-error {{missing map type modifier}} expected-error {{missing map type modifier}} expected-error {{missing map type}}
     {}
-    #pragma omp target map( d, f, bf: a)   // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
+    #pragma omp target map( d, f, bf: a)   // expected-error 2 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
     {}
-    #pragma omp target map( , f, : a)   // expected-error {{missing map type modifier}} expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+    #pragma omp target map( , f, : a)   // expected-error {{missing map type modifier}} expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper'}} expected-error {{missing map type}}
     {}
     #pragma omp target map(always close: a)   // expected-error {{missing map type}}
     {}
     #pragma omp target map(always close bf: a)   // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
     {}
-    #pragma omp target map(always tofrom close: a)   // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+    #pragma omp target map(always tofrom close: a)   // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
     {}
-    #pragma omp target map(tofrom from: a)   // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+    #pragma omp target map(tofrom from: a)   // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
     {}
     #pragma omp target map(close bf: a)   // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
     {}
@@ -467,7 +479,7 @@
   T *k = &j;
   T x;
   T y;
-  T to, tofrom, always, close;
+  T to, tofrom, always, close, present;
   const T (&l)[5] = da;
 #pragma omp target map // expected-error {{expected '(' after 'map'}}
   {}
@@ -540,17 +552,23 @@
 
 #pragma omp target data map(always, tofrom: x)
 #pragma omp target data map(always: x) // expected-error {{missing map type}}
-#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
 #pragma omp target data map(always, tofrom: always, tofrom, x)
 #pragma omp target map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   foo();
 
 #pragma omp target data map(close, tofrom: x)
 #pragma omp target data map(close: x) // expected-error {{missing map type}}
-#pragma omp target data map(tofrom, close: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target data map(tofrom, close: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
 #pragma omp target data map(close, tofrom: close, tofrom, x)
   foo();
 
+#pragma omp target data map(present, tofrom: x)
+#pragma omp target data map(present: x) // expected-error {{missing map type}}
+#pragma omp target data map(tofrom, present: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
+#pragma omp target data map(present, tofrom: present, tofrom, x)
+  foo();
+
   T marr[10][10], iarr[5];
 #pragma omp target data map(marr[10][0:2:2]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
   {}
@@ -600,7 +618,7 @@
   S6<int> m;
   int x;
   int y;
-  int to, tofrom, always, close;
+  int to, tofrom, always, close, present;
   const int (&l)[5] = da;
   SC1 s;
   SC1 *p;
@@ -654,13 +672,17 @@
 
 #pragma omp target data map(always, tofrom: x)
 #pragma omp target data map(always: x) // expected-error {{missing map type}}
-#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
 #pragma omp target data map(always, tofrom: always, tofrom, x)
 #pragma omp target map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   foo();
 #pragma omp target data map(close, tofrom: x)
 #pragma omp target data map(close: x) // expected-error {{missing map type}}
-#pragma omp target data map(tofrom, close: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target data map(tofrom, close: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
+  foo();
+#pragma omp target data map(present, tofrom: x)
+#pragma omp target data map(present: x) // expected-error {{missing map type}}
+#pragma omp target data map(tofrom, present: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
   foo();
 #pragma omp target private(j) map(j) // expected-error {{private variable cannot be in a map clause in '#pragma omp target' directive}}  expected-note {{defined as private}}
   {}
Index: clang/test/OpenMP/target_parallel_for_map_messages.cpp
===================================================================
--- clang/test/OpenMP/target_parallel_for_map_messages.cpp
+++ clang/test/OpenMP/target_parallel_for_map_messages.cpp
@@ -175,7 +175,7 @@
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
@@ -287,7 +287,7 @@
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
Index: clang/test/OpenMP/target_parallel_for_simd_map_messages.cpp
===================================================================
--- clang/test/OpenMP/target_parallel_for_simd_map_messages.cpp
+++ clang/test/OpenMP/target_parallel_for_simd_map_messages.cpp
@@ -175,7 +175,7 @@
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
@@ -287,7 +287,7 @@
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
Index: clang/test/OpenMP/target_parallel_map_messages.cpp
===================================================================
--- clang/test/OpenMP/target_parallel_map_messages.cpp
+++ clang/test/OpenMP/target_parallel_map_messages.cpp
@@ -175,7 +175,7 @@
   foo();
 #pragma omp target parallel map(always: x) // expected-error {{missing map type}}
   foo();
-#pragma omp target parallel map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target parallel map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
   foo();
 #pragma omp target parallel map(always, tofrom: always, tofrom, x)
   foo();
@@ -285,7 +285,7 @@
   foo();
 #pragma omp target parallel map(always: x) // expected-error {{missing map type}}
   foo();
-#pragma omp target parallel map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target parallel map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
   foo();
 #pragma omp target parallel map(always, tofrom: always, tofrom, x)
   foo();
Index: clang/test/OpenMP/target_simd_map_messages.cpp
===================================================================
--- clang/test/OpenMP/target_simd_map_messages.cpp
+++ clang/test/OpenMP/target_simd_map_messages.cpp
@@ -169,7 +169,7 @@
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
@@ -275,7 +275,7 @@
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
Index: clang/test/OpenMP/target_teams_distribute_map_messages.cpp
===================================================================
--- clang/test/OpenMP/target_teams_distribute_map_messages.cpp
+++ clang/test/OpenMP/target_teams_distribute_map_messages.cpp
@@ -175,7 +175,7 @@
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target teams distribute map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
@@ -287,7 +287,7 @@
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target teams distribute map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
Index: clang/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp
===================================================================
--- clang/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp
+++ clang/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp
@@ -174,7 +174,7 @@
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target teams distribute parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
@@ -285,7 +285,7 @@
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target teams distribute parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
Index: clang/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp
===================================================================
--- clang/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp
+++ clang/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp
@@ -175,7 +175,7 @@
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target teams distribute parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
@@ -286,7 +286,7 @@
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target teams distribute parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
Index: clang/test/OpenMP/target_teams_distribute_simd_map_messages.cpp
===================================================================
--- clang/test/OpenMP/target_teams_distribute_simd_map_messages.cpp
+++ clang/test/OpenMP/target_teams_distribute_simd_map_messages.cpp
@@ -175,7 +175,7 @@
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target teams distribute simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
@@ -287,7 +287,7 @@
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target teams distribute simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
Index: clang/test/OpenMP/target_teams_map_messages.cpp
===================================================================
--- clang/test/OpenMP/target_teams_map_messages.cpp
+++ clang/test/OpenMP/target_teams_map_messages.cpp
@@ -468,7 +468,7 @@
 
 #pragma omp target data map(always, tofrom: x)
 #pragma omp target data map(always: x) // expected-error {{missing map type}}
-#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
 #pragma omp target data map(always, tofrom: always, tofrom, x)
 #pragma omp target teams map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   foo();
@@ -543,7 +543,7 @@
 
 #pragma omp target data map(always, tofrom: x)
 #pragma omp target data map(always: x) // expected-error {{missing map type}}
-#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}} expected-error {{missing map type}}
 #pragma omp target data map(always, tofrom: always, tofrom, x)
 #pragma omp target teams map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   foo();