diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7384,10 +7384,9 @@
     // &p, &p, sizeof(float*), TARGET_PARAM | TO | FROM
     //
     // map(p[1:24])
+    // &p, &p[1], 24*sizeof(float), TARGET_PARAM | TO | FROM | PTR_AND_OBJ
+    // in unified shared memory mode
     // p, &p[1], 24*sizeof(float), TARGET_PARAM | TO | FROM
-    // for data directives
-    // p, p, sizeof(float*), TARGET_PARAM | TO | FROM
-    // p, &p[1], 24*sizeof(float), PTR_AND_OBJ | TO | FROM
     //
     // map(s)
     // &s, &s, sizeof(S2), TARGET_PARAM | TO | FROM
@@ -7522,6 +7521,7 @@
     // Track if the map information being generated is the first for a list of
     // components.
     bool IsExpressionFirstInfo = true;
+    bool FirstPointerInComplexData = false;
     Address BP = Address::invalid();
     const Expr *AssocExpr = I->getAssociatedExpression();
     const auto *AE = dyn_cast<ArraySubscriptExpr>(AssocExpr);
@@ -7564,17 +7564,14 @@
       QualType Ty =
           I->getAssociatedDeclaration()->getType().getNonReferenceType();
       if (Ty->isAnyPointerType() && std::next(I) != CE) {
-        BP = CGF.EmitLoadOfPointer(BP, Ty->castAs<PointerType>());
-
-        // For non-data directives, we do not need to generate individual map
-        // information for the pointer, it can be associated with the combined
-        // storage.
-        if (CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory() ||
-            !CurDir.is<const OMPExecutableDirective *>() ||
-            !isOpenMPTargetDataManagementDirective(
-                CurDir.get<const OMPExecutableDirective *>()
-                    ->getDirectiveKind()))
-          ++I;
+        // No need to generate individual map information for the pointer, it
+        // can be associated with the combined storage if shared memory mode is
+        // active.
+        if (CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory())
+          BP = CGF.EmitLoadOfPointer(BP, Ty->castAs<PointerType>());
+        else
+          FirstPointerInComplexData = true;
+        ++I;
       }
     }
 
@@ -7609,8 +7606,19 @@
         EncounteredME = dyn_cast<MemberExpr>(I->getAssociatedExpression());
         // If we encounter a PTR_AND_OBJ entry from now on it should be marked
         // as MEMBER_OF the parent struct.
-        if (EncounteredME)
+        if (EncounteredME) {
           ShouldBeMemberOf = true;
+          // Do not emit as complex pointer if this is actually not array-like
+          // expression.
+          if (FirstPointerInComplexData) {
+            QualType Ty = std::prev(I)
+                              ->getAssociatedDeclaration()
+                              ->getType()
+                              .getNonReferenceType();
+            BP = CGF.EmitLoadOfPointer(BP, Ty->castAs<PointerType>());
+            FirstPointerInComplexData = false;
+          }
+        }
       }
 
       auto Next = std::next(I);
@@ -7750,10 +7758,11 @@
           // same expression except for the first one. We also need to signal
           // this map is the first one that relates with the current capture
           // (there is a set of entries for each capture).
-          OpenMPOffloadMappingFlags Flags = getMapTypeBits(
-              MapType, MapModifiers, IsImplicit,
-              !IsExpressionFirstInfo || RequiresReference,
-              IsCaptureFirstInfo && !RequiresReference);
+          OpenMPOffloadMappingFlags Flags =
+              getMapTypeBits(MapType, MapModifiers, IsImplicit,
+                             !IsExpressionFirstInfo || RequiresReference ||
+                                 FirstPointerInComplexData,
+                             IsCaptureFirstInfo && !RequiresReference);
 
           if (!IsExpressionFirstInfo) {
             // If we have a PTR_AND_OBJ pair where the OBJ is a pointer as well,
@@ -7811,6 +7820,7 @@
 
         IsExpressionFirstInfo = false;
         IsCaptureFirstInfo = false;
+        FirstPointerInComplexData = false;
       }
     }
   }
@@ -8055,6 +8065,7 @@
     // emission of that entry until the whole struct has been processed.
     llvm::MapVector<const ValueDecl *, SmallVector<DeferredDevicePtrEntryTy, 4>>
         DeferredInfo;
+    MapCombinedInfoTy UseDevicePtrCombinedInfo;
 
     for (const auto *C :
          CurExecDir->getClausesOfKind<OMPUseDevicePtrClause>()) {
@@ -8074,13 +8085,22 @@
         // We potentially have map information for this declaration already.
         // Look for the first set of components that refer to it.
         if (It != Info.end()) {
-          auto CI = std::find_if(
-              It->second.begin(), It->second.end(), [VD](const MapInfo &MI) {
-                return MI.Components.back().getAssociatedDeclaration() == VD;
-              });
+          auto *CI = llvm::find_if(It->second, [VD](const MapInfo &MI) {
+            return MI.Components.back().getAssociatedDeclaration() == VD;
+          });
           // If we found a map entry, signal that the pointer has to be returned
           // and move on to the next declaration.
-          if (CI != It->second.end()) {
+          // Exclude cases where the base pointer is mapped as array subscript,
+          // array section or array shaping. The base address is passed as a
+          // pointer to base in this case and cannot be used as a base for
+          // use_device_ptr list item.
+          auto PrevCI = std::next(CI->Components.rbegin());
+          if (CI != It->second.end() &&
+              (CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory() ||
+               isa<MemberExpr>(IE) ||
+               !VD->getType().getNonReferenceType()->isPointerType() ||
+               PrevCI == CI->Components.rend() ||
+               isa<MemberExpr>(PrevCI->getAssociatedExpression()))) {
             CI->ReturnDevicePointer = true;
             continue;
           }
@@ -8103,13 +8123,13 @@
         } else {
           llvm::Value *Ptr =
               CGF.EmitLoadOfScalar(CGF.EmitLValue(IE), IE->getExprLoc());
-          CombinedInfo.BasePointers.emplace_back(Ptr, VD);
-          CombinedInfo.Pointers.push_back(Ptr);
-          CombinedInfo.Sizes.push_back(
+          UseDevicePtrCombinedInfo.BasePointers.emplace_back(Ptr, VD);
+          UseDevicePtrCombinedInfo.Pointers.push_back(Ptr);
+          UseDevicePtrCombinedInfo.Sizes.push_back(
               llvm::Constant::getNullValue(CGF.Int64Ty));
-          CombinedInfo.Types.push_back(OMP_MAP_RETURN_PARAM |
-                                       OMP_MAP_TARGET_PARAM);
-          CombinedInfo.Mappers.push_back(nullptr);
+          UseDevicePtrCombinedInfo.Types.push_back(OMP_MAP_RETURN_PARAM |
+                                                   OMP_MAP_TARGET_PARAM);
+          UseDevicePtrCombinedInfo.Mappers.push_back(nullptr);
         }
       }
     }
@@ -8260,6 +8280,8 @@
       // We need to append the results of this capture to what we already have.
       CombinedInfo.append(CurInfo);
     }
+    // Append data for use_device_ptr clauses.
+    CombinedInfo.append(UseDevicePtrCombinedInfo);
   }
 
   /// Generate all the base pointers, section pointers, sizes, map types, and
diff --git a/clang/test/OpenMP/target_data_codegen.cpp b/clang/test/OpenMP/target_data_codegen.cpp
--- a/clang/test/OpenMP/target_data_codegen.cpp
+++ b/clang/test/OpenMP/target_data_codegen.cpp
@@ -555,7 +555,7 @@
 
 void test_close_modifier(int arg) {
   S2 *ps;
-  // CK5: private unnamed_addr constant [6 x i64] [i64 1059, i64 32, i64 562949953422339, i64 562949953421328, i64 16, i64 1043]
+  // CK5: private unnamed_addr constant [5 x i64] [i64 1059, i64 32, i64 562949953421328, i64 16, i64 1043]
   #pragma omp target data map(close,tofrom: arg, ps->ps->ps->ps->s)
   {
     ++(arg);
@@ -634,20 +634,17 @@
   // Make sure the struct picks up present even if another element of the struct
   // doesn't have present.
 
-  // CK8: private unnamed_addr constant [15 x i64]
+  // CK8: private unnamed_addr constant [11 x i64]
 
   // ps1
   //
   // PRESENT=0x1000 | TARGET_PARAM=0x20 = 0x1020
   // MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003
-  // MEMBER_OF_1=0x1000000000000 | PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x1000000000013
-  // MEMBER_OF_1=0x1000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x1000000001003
   // MEMBER_OF_1=0x1000000000000 | PRESENT=0x1000 | PTR_AND_OBJ=0x10 = 0x1000000001010
   // PRESENT=0x1000 | PTR_AND_OBJ=0x10 = 0x1010
   // PRESENT=0x1000 | PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x1013
   //
   // CK8-SAME: {{^}} [i64 [[#0x1020]], i64 [[#0x1000000000003]],
-  // CK8-SAME: {{^}} i64 [[#0x1000000000013]], i64 [[#0x1000000001003]],
   // CK8-SAME: {{^}} i64 [[#0x1000000001010]], i64 [[#0x1010]], i64 [[#0x1013]],
 
   // arg
@@ -659,16 +656,13 @@
   // ps2
   //
   // PRESENT=0x1000 | TARGET_PARAM=0x20 = 0x1020
-  // MEMBER_OF_9=0x9000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x9000000001003
-  // MEMBER_OF_9=0x9000000000000 | PRESENT=0x1000 | PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x9000000001013
-  // MEMBER_OF_9=0x9000000000000 | FROM=0x2 | TO=0x1 = 0x9000000000003
-  // MEMBER_OF_9=0x9000000000000 | PTR_AND_OBJ=0x10 = 0x9000000000010
+  // MEMBER_OF_7=0x9000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x7000000001003
+  // MEMBER_OF_7=0x9000000000000 | PTR_AND_OBJ=0x10 = 0x7000000000010
   // PTR_AND_OBJ=0x10 = 0x10
   // PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x13
   //
-  // CK8-SAME: {{^}} i64 [[#0x1020]], i64 [[#0x9000000001003]],
-  // CK8-SAME: {{^}} i64 [[#0x9000000001013]], i64 [[#0x9000000000003]],
-  // CK8-SAME: {{^}} i64 [[#0x9000000000010]], i64 [[#0x10]], i64 [[#0x13]]]
+  // CK8-SAME: {{^}} i64 [[#0x1020]], i64 [[#0x7000000001003]],
+  // CK8-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)
diff --git a/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp b/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp
--- a/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp
+++ b/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp
@@ -22,18 +22,18 @@
 double *g;
 
 // CK1: @g = global double*
-// CK1: [[MTYPE00:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 19]
-// CK1: [[MTYPE01:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 19]
-// CK1: [[MTYPE03:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 19]
-// CK1: [[MTYPE04:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 19]
-// CK1: [[MTYPE05:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 19]
-// CK1: [[MTYPE06:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 19]
-// CK1: [[MTYPE07:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 19]
-// CK1: [[MTYPE08:@.+]] = {{.*}}constant [4 x i64] [i64 99, i64 19, i64 35, i64 19]
-// CK1: [[MTYPE09:@.+]] = {{.*}}constant [4 x i64] [i64 99, i64 19, i64 99, i64 19]
-// CK1: [[MTYPE10:@.+]] = {{.*}}constant [4 x i64] [i64 99, i64 19, i64 99, i64 19]
-// CK1: [[MTYPE11:@.+]] = {{.*}}constant [3 x i64] [i64 96, i64 35, i64 19]
-// CK1: [[MTYPE12:@.+]] = {{.*}}constant [3 x i64] [i64 96, i64 35, i64 19]
+// CK1: [[MTYPE00:@.+]] = {{.*}}constant [2 x i64] [i64 51, i64 96]
+// CK1: [[MTYPE01:@.+]] = {{.*}}constant [2 x i64] [i64 51, i64 96]
+// CK1: [[MTYPE03:@.+]] = {{.*}}constant [2 x i64] [i64 51, i64 96]
+// CK1: [[MTYPE04:@.+]] = {{.*}}constant [2 x i64] [i64 51, i64 96]
+// CK1: [[MTYPE05:@.+]] = {{.*}}constant [2 x i64] [i64 51, i64 96]
+// CK1: [[MTYPE06:@.+]] = {{.*}}constant [2 x i64] [i64 51, i64 96]
+// CK1: [[MTYPE07:@.+]] = {{.*}}constant [2 x i64] [i64 51, i64 96]
+// CK1: [[MTYPE08:@.+]] = {{.*}}constant [3 x i64] [i64 51, i64 51, i64 96]
+// CK1: [[MTYPE09:@.+]] = {{.*}}constant [4 x i64] [i64 51, i64 51, i64 96, i64 96]
+// CK1: [[MTYPE10:@.+]] = {{.*}}constant [4 x i64] [i64 51, i64 51, i64 96, i64 96]
+// CK1: [[MTYPE11:@.+]] = {{.*}}constant [2 x i64] [i64 51, i64 96]
+// CK1: [[MTYPE12:@.+]] = {{.*}}constant [2 x i64] [i64 51, i64 96]
 
 // CK1-LABEL: @_Z3foo
 template<typename T>
@@ -41,13 +41,14 @@
   float *l;
   T *t;
 
-  // CK1:     [[T:%.+]] = load double*, double** [[DECL:@g]],
   // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0
+  // CK1:     [[CBP:%.+]] = bitcast i8** [[BP]] to double***
+  // CK1:     store double** @g, double*** [[CBP]],
+  // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 1
   // CK1:     [[CBP:%.+]] = bitcast i8** [[BP]] to double**
-  // CK1:     store double* [[T]], double** [[CBP]],
   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE00]]
   // CK1:     [[VAL:%.+]] = load double*, double** [[CBP]],
-  // CK1-NOT: store double* [[VAL]], double** [[DECL]],
+  // CK1-NOT: store double* [[VAL]], double** @g,
   // CK1:     store double* [[VAL]], double** [[PVT:%.+]],
   // CK1:     [[TT:%.+]] = load double*, double** [[PVT]],
   // CK1:     getelementptr inbounds double, double* [[TT]], i32 1
@@ -56,14 +57,15 @@
     ++g;
   }
   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE00]]
-  // CK1:     [[TTT:%.+]] = load double*, double** [[DECL]],
+  // CK1:     [[TTT:%.+]] = load double*, double** @g,
   // CK1:     getelementptr inbounds double, double* [[TTT]], i32 1
   ++g;
 
-  // CK1:     [[T1:%.+]] = load float*, float** [[DECL:%.+]],
   // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0
-  // CK1:     [[CBP:%.+]] = bitcast i8** [[BP]] to float**
-  // CK1:     store float* [[T1]], float** [[CBP]],
+  // CK1:     [[CBP:%.+]] = bitcast i8** [[BP]] to float***
+  // CK1:     store float** [[DECL:%.+]], float*** [[CBP]],
+  // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 1
+  // CK1:     [[CBP:%.+]] = bitcast i8** [[BP]] to float*
   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE01]]
   // CK1:     [[VAL:%.+]] = load float*, float** [[CBP]],
   // CK1-NOT: store float* [[VAL]], float** [[DECL]],
@@ -91,10 +93,11 @@
   // CK1:     getelementptr inbounds float, float* [[TTT]], i32 1
   ++l;
 
-  // CK1:     [[T1:%.+]] = load float*, float** [[DECL:%.+]],
   // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0
+  // CK1:     [[CBP:%.+]] = bitcast i8** [[BP]] to float***
+  // CK1:     store float** [[DECL:%.+]], float*** [[CBP]],
+  // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 1
   // CK1:     [[CBP:%.+]] = bitcast i8** [[BP]] to float**
-  // CK1:     store float* [[T1]], float** [[CBP]],
   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE03]]
   // CK1:     [[VAL:%.+]] = load float*, float** [[CBP]],
   // CK1-NOT: store float* [[VAL]], float** [[DECL]],
@@ -114,10 +117,11 @@
   // CK1:     br i1 [[CMP]], label %[[BTHEN:.+]], label %[[BELSE:.+]]
 
   // CK1:     [[BTHEN]]:
-  // CK1:     [[T1:%.+]] = load float*, float** [[DECL:%.+]],
   // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0
+  // CK1:     [[CBP:%.+]] = bitcast i8** [[BP]] to float***
+  // CK1:     store float** [[DECL:%.+]], float*** [[CBP]],
+  // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 1
   // CK1:     [[CBP:%.+]] = bitcast i8** [[BP]] to float**
-  // CK1:     store float* [[T1]], float** [[CBP]],
   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE04]]
   // CK1:     [[VAL:%.+]] = load float*, float** [[CBP]],
   // CK1-NOT: store float* [[VAL]], float** [[DECL]],
@@ -151,10 +155,11 @@
   ++l;
 
   // CK1:     [[T2:%.+]] = load float**, float*** [[DECL:%.+]],
-  // CK1:     [[T1:%.+]] = load float*, float** [[T2]],
   // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0
+  // CK1:     [[CBP:%.+]] = bitcast i8** [[BP]] to float***
+  // CK1:     store float** [[T2:%.+]], float*** [[CBP]],
+  // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 1
   // CK1:     [[CBP:%.+]] = bitcast i8** [[BP]] to float**
-  // CK1:     store float* [[T1]], float** [[CBP]],
   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE05]]
   // CK1:     [[VAL:%.+]] = load float*, float** [[CBP]],
   // CK1:     store float* [[VAL]], float** [[PVTV:%.+]],
@@ -173,10 +178,11 @@
   // CK1:     getelementptr inbounds float, float* [[TTTT]], i32 1
   ++lr;
 
-  // CK1:     [[T1:%.+]] = load i32*, i32** [[DECL:%.+]],
   // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0
+  // CK1:     [[CBP:%.+]] = bitcast i8** [[BP]] to i32***
+  // CK1:     store i32** [[DECL:%.+]], i32*** [[CBP]],
+  // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 1
   // CK1:     [[CBP:%.+]] = bitcast i8** [[BP]] to i32**
-  // CK1:     store i32* [[T1]], i32** [[CBP]],
   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE06]]
   // CK1:     [[VAL:%.+]] = load i32*, i32** [[CBP]],
   // CK1-NOT: store i32* [[VAL]], i32** [[DECL]],
@@ -192,11 +198,13 @@
   // CK1:     getelementptr inbounds i32, i32* [[TTT]], i32 1
   ++t;
 
+  // CK1:     load i32**, i32*** %
   // CK1:     [[T2:%.+]] = load i32**, i32*** [[DECL:%.+]],
-  // CK1:     [[T1:%.+]] = load i32*, i32** [[T2]],
   // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0
+  // CK1:     [[CBP:%.+]] = bitcast i8** [[BP]] to i32***
+  // CK1:     store i32** [[T2]], i32*** [[CBP]],
+  // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 1
   // CK1:     [[CBP:%.+]] = bitcast i8** [[BP]] to i32**
-  // CK1:     store i32* [[T1]], i32** [[CBP]],
   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE07]]
   // CK1:     [[VAL:%.+]] = load i32*, i32** [[CBP]],
   // CK1:     store i32* [[VAL]], i32** [[PVTV:%.+]],
@@ -215,10 +223,11 @@
   // CK1:     getelementptr inbounds i32, i32* [[TTTT]], i32 1
   ++tr;
 
-  // CK1:     [[T1:%.+]] = load float*, float** [[DECL:%.+]],
-  // CK1:     [[BP:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* %{{.+}}, i32 0, i32
+  // CK1:     [[BP:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %{{.+}}, i32 0, i32 0
+  // CK1:     [[CBP:%.+]] = bitcast i8** [[BP]] to float***
+  // CK1:     store float** [[DECL:%.+]], float*** [[CBP]],
+  // CK1:     [[BP:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %{{.+}}, i32 0, i32 2
   // CK1:     [[CBP:%.+]] = bitcast i8** [[BP]] to float**
-  // CK1:     store float* [[T1]], float** [[CBP]],
   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE08]]
   // CK1:     [[VAL:%.+]] = load float*, float** [[CBP]],
   // CK1-NOT: store float* [[VAL]], float** [[DECL]],
@@ -235,6 +244,8 @@
   ++l; ++t;
 
 
+  // CK1:     bitcast i8** {{%.+}} to float***
+  // CK1:     bitcast i8** {{%.+}} to i32***
   // CK1:     [[_CBP:%.+]] = bitcast i8** {{%.+}} to float**
   // CK1:     [[CBP:%.+]] = bitcast i8** {{%.+}} to i32**
   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE09]]
@@ -257,6 +268,8 @@
   // CK1:     getelementptr inbounds i32, i32* [[TTT]], i32 1
   ++l; ++t;
 
+  // CK1:     bitcast i8** {{%.+}} to float***
+  // CK1:     bitcast i8** {{%.+}} to i32***
   // CK1:     [[_CBP:%.+]] = bitcast i8** {{%.+}} to float**
   // CK1:     [[CBP:%.+]] = bitcast i8** {{%.+}} to i32**
   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE10]]
@@ -280,7 +293,7 @@
   ++l; ++t;
 
   // CK1:     [[T1:%.+]] = load i32*, i32** [[DECL:%.+]],
-  // CK1:     [[BP:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %{{.+}}, i32 0, i32 0
+  // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 1
   // CK1:     [[CBP:%.+]] = bitcast i8** [[BP]] to i32**
   // CK1:     store i32* [[T1]], i32** [[CBP]],
   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE11]]
@@ -300,7 +313,7 @@
 
   // CK1:     [[T2:%.+]] = load i32**, i32*** [[DECL:%.+]],
   // CK1:     [[T1:%.+]] = load i32*, i32** [[T2]],
-  // CK1:     [[BP:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %{{.+}}, i32 0, i32 0
+  // CK1:     [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 1
   // CK1:     [[CBP:%.+]] = bitcast i8** [[BP]] to i32**
   // CK1:     store i32* [[T1]], i32** [[CBP]],
   // CK1:     call void @__tgt_target_data_begin{{.+}}[[MTYPE12]]
@@ -348,7 +361,7 @@
 // CK2: [[ST:%.+]] = type { double*, double** }
 // CK2: [[MTYPE00:@.+]] = {{.*}}constant [2 x i64] [i64 32, i64 281474976710739]
 // CK2: [[MTYPE01:@.+]] = {{.*}}constant [2 x i64] [i64 32, i64 281474976710739]
-// CK2: [[MTYPE02:@.+]] = {{.*}}constant [4 x i64] [i64 35, i64 19, i64 32, i64 844424930132048]
+// CK2: [[MTYPE02:@.+]] = {{.*}}constant [3 x i64] [i64 51, i64 32, i64 562949953421392]
 // CK2: [[MTYPE03:@.+]] = {{.*}}constant [3 x i64] [i64 32, i64 281474976710739, i64 281474976710736]
 
 template <typename T>
@@ -404,7 +417,7 @@
     // CK2:     getelementptr inbounds double, double* [[TTTT]], i32 1
     b++;
 
-    // CK2:     [[BP:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* %{{.+}}, i32 0, i32 3
+    // CK2:     [[BP:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %{{.+}}, i32 0, i32 2
     // CK2:     [[CBP:%.+]] = bitcast i8** [[BP]] to double***
     // CK2:     store double** [[RVAL:%.+]], double*** [[CBP]],
     // CK2:     call void @__tgt_target_data_begin{{.+}}[[MTYPE02]]
diff --git a/clang/test/OpenMP/target_map_codegen.cpp b/clang/test/OpenMP/target_map_codegen.cpp
--- a/clang/test/OpenMP/target_map_codegen.cpp
+++ b/clang/test/OpenMP/target_map_codegen.cpp
@@ -1657,25 +1657,25 @@
 
 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK19: [[SIZE10:@.+]] = private {{.*}}constant [1 x i64] [i64 240]
-// CK19: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
+// CK19: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i64] [i64 51]
 
 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK19: [[SIZE11:@.+]] = private {{.*}}constant [1 x i64] [i64 240]
-// CK19: [[MTYPE11:@.+]] = private {{.*}}constant [1 x i64] [i64 32]
+// CK19: [[MTYPE11:@.+]] = private {{.*}}constant [1 x i64] [i64 48]
 
 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK19: [[SIZE12:@.+]] = private {{.*}}constant [1 x i64] [i64 4]
-// CK19: [[MTYPE12:@.+]] = private {{.*}}constant [1 x i64] [i64 33]
+// CK19: [[MTYPE12:@.+]] = private {{.*}}constant [1 x i64] [i64 49]
 
 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
-// CK19: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i64] [i64 32]
+// CK19: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i64] [i64 48]
 
 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
-// CK19: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i64] [i64 33]
+// CK19: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i64] [i64 49]
 
 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK19: [[SIZE15:@.+]] = private {{.*}}constant [1 x i64] [i64 4]
-// CK19: [[MTYPE15:@.+]] = private {{.*}}constant [1 x i64] [i64 34]
+// CK19: [[MTYPE15:@.+]] = private {{.*}}constant [1 x i64] [i64 50]
 
 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK19-USE: [[MTYPE16:@.+]] = private {{.*}}constant [2 x i64] [i64 800, i64 33]
@@ -1735,11 +1735,11 @@
 
 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK19: [[SIZE28:@.+]] = private {{.*}}constant [3 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 16]
-// CK19: [[MTYPE28:@.+]] = private {{.*}}constant [3 x i64] [i64 35, i64 16, i64 19]
+// CK19: [[MTYPE28:@.+]] = private {{.*}}constant [3 x i64] [i64 51, i64 16, i64 19]
 
 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK19: [[SIZE29:@.+]] = private {{.*}}constant [3 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 4]
-// CK19: [[MTYPE29:@.+]] = private {{.*}}constant [3 x i64] [i64 35, i64 16, i64 19]
+// CK19: [[MTYPE29:@.+]] = private {{.*}}constant [3 x i64] [i64 51, i64 16, i64 19]
 
 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK19-USE: [[MTYPE30:@.+]] = private {{.*}}constant [4 x i64] [i64 800, i64 800, i64 800, i64 35]
@@ -1794,7 +1794,7 @@
 
 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK19: [[SIZE42:@.+]] = private {{.*}}constant [3 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 104]
-// CK19: [[MTYPE42:@.+]] = private {{.*}}constant [3 x i64] [i64 35, i64 16, i64 19]
+// CK19: [[MTYPE42:@.+]] = private {{.*}}constant [3 x i64] [i64 51, i64 16, i64 19]
 
 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK19: [[MTYPE43:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
@@ -2071,11 +2071,10 @@
 
   // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-  // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**
+  // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32***
   // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
-  // CK19-DAG: store i32* [[RVAR0:%.+]], i32** [[CBP0]]
+  // CK19-DAG: store i32** [[VAR0:%.+]], i32*** [[CBP0]]
   // CK19-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]]
-  // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]]
   // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 20
   // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]]
 
@@ -2095,11 +2094,10 @@
 
   // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-  // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**
+  // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32***
   // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
-  // CK19-DAG: store i32* [[RVAR0:%.+]], i32** [[CBP0]]
+  // CK19-DAG: store i32** [[VAR0:%.+]], i32*** [[CBP0]]
   // CK19-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]]
-  // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]]
   // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 0
   // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]]
 
@@ -2119,11 +2117,10 @@
 
   // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-  // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**
+  // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32***
   // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
-  // CK19-DAG: store i32* [[RVAR0:%.+]], i32** [[CBP0]]
+  // CK19-DAG: store i32** [[VAR0:%.+]], i32*** [[CBP0]]
   // CK19-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]]
-  // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]]
   // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 15
   // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]]
 
@@ -2145,13 +2142,12 @@
   // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
   // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
-  // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**
+  // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32***
   // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
-  // CK19-DAG: store i32* [[RVAR0:%.+]], i32** [[CBP0]]
+  // CK19-DAG: store i32** [[VAR0:%.+]], i32*** [[CBP0]]
   // CK19-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]]
   // CK19-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], i{{.+}}* [[S0]]
   // CK19-DAG: [[CSVAL0]] = {{mul nuw i64 %.*, 4|sext i32 .+ to i64}}
-  // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]]
   // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} %{{.*}}
   // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]]
 
@@ -2173,13 +2169,12 @@
   // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
   // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
-  // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**
+  // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32***
   // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
-  // CK19-DAG: store i32* [[RVAR0:%.+]], i32** [[CBP0]]
+  // CK19-DAG: store i32** [[VAR0:%.+]], i32*** [[CBP0]]
   // CK19-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]]
   // CK19-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], i{{.+}}* [[S0]]
   // CK19-DAG: [[CSVAL0]] = {{mul nuw i64 %.*, 4|sext i32 .+ to i64}}
-  // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]]
   // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 0
   // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]]
 
@@ -2199,11 +2194,10 @@
 
   // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-  // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**
+  // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32***
   // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
-  // CK19-DAG: store i32* [[RVAR0:%.+]], i32** [[CBP0]]
+  // CK19-DAG: store i32** [[VAR0:%.+]], i32*** [[CBP0]]
   // CK19-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]]
-  // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]]
   // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} %{{.*}}
   // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]]
 
@@ -2629,11 +2623,10 @@
 
   // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-  // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32****
+  // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32*****
   // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32****
-  // CK19-DAG: store i32*** [[VAR0:%.+]], i32**** [[CBP0]]
+  // CK19-DAG: store i32**** [[PTR:%.+]], i32***** [[CBP0]]
   // CK19-DAG: store i32*** [[SEC0:%.+]], i32**** [[CP0]]
-  // CK19-DAG: [[VAR0]] = load i32***, i32**** [[PTR:%[^,]+]],
   // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32*** [[SEC00:[^,]+]], i{{.+}} 1
   // CK19-DAG: [[SEC00]] = load i32***, i32**** [[PTR]],
 
@@ -2677,11 +2670,10 @@
 
   // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-  // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32****
+  // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32*****
   // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32****
-  // CK19-DAG: store i32*** [[VAR0:%.+]], i32**** [[CBP0]]
+  // CK19-DAG: store i32**** [[PTR:%.+]], i32***** [[CBP0]]
   // CK19-DAG: store i32*** [[SEC0:%.+]], i32**** [[CP0]]
-  // CK19-DAG: [[VAR0]] = load i32***, i32**** [[PTR:%[^,]+]],
   // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32*** [[SEC00:[^,]+]], i{{.+}} 1
   // CK19-DAG: [[SEC00]] = load i32***, i32**** [[PTR]],
 
@@ -3247,11 +3239,10 @@
 
   // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-  // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to double****
+  // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to double*****
   // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to double****
-  // CK19-DAG: store double*** [[VAR0:%.+]], double**** [[CBP0]]
+  // CK19-DAG: store double**** [[PTR:%.+]], double***** [[CBP0]]
   // CK19-DAG: store double*** [[SEC0:%.+]], double**** [[CP0]]
-  // CK19-DAG: [[VAR0]] = load double***, double**** [[PTR:%[^,]+]],
   // CK19-DAG: [[SEC0]] = getelementptr {{.*}}double*** [[SEC00:[^,]+]], i{{.+}} 0
   // CK19-DAG: [[SEC00]] = load double***, double**** [[PTR]],
 
@@ -3432,7 +3423,7 @@
 
 // CK20-LABEL: @.__omp_offloading_{{.*}}explicit_maps_references_and_function_args{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK20: [[SIZE03:@.+]] = private {{.*}}constant [1 x i64] [i64 12]
-// CK20: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i64] [i64 34]
+// CK20: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i64] [i64 50]
 
 // CK20-LABEL: explicit_maps_references_and_function_args{{.*}}(
 void explicit_maps_references_and_function_args (int a, float b, int (&c)[10], float *d){
@@ -3508,11 +3499,10 @@
 
   // CK20-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK20-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-  // CK20-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
+  // CK20-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float***
   // CK20-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
-  // CK20-DAG: store float* [[RVAR0:%.+]], float** [[CBP0]]
+  // CK20-DAG: store float** [[VAR0:%.+]], float*** [[CBP0]]
   // CK20-DAG: store float* [[SEC0:%.+]], float** [[CP0]]
-  // CK20-DAG: [[RVAR0]] = load float*, float** [[VAR0:%[^,]+]]
   // CK20-DAG: [[SEC0]] = getelementptr {{.*}}float* [[RVAR00:%.+]], i{{.+}} 2
   // CK20-DAG: [[RVAR00]] = load float*, float** [[VAR0]]
 
@@ -3595,7 +3585,7 @@
 
 // CK21-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK21: [[SIZE01:@.+]] = private {{.*}}constant [1 x i64] [i64 492]
-// CK21: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
+// CK21: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 51]
 
 // CK21-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK21: [[MTYPE02:@.+]] = private {{.*}}constant [2 x i64] [i64 32, i64 281474976710674]
@@ -3665,11 +3655,10 @@
 
     // CK21-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
     // CK21-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-    // CK21-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**
+    // CK21-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32***
     // CK21-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
-    // CK21-DAG: store i32* [[RVAR0:%.+]], i32** [[CBP0]]
+    // CK21-DAG: store i32** [[VAR0:%.+]], i32*** [[CBP0]]
     // CK21-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]]
-    // CK21-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]]
     // CK21-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 0
     // CK21-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]]
 
@@ -3874,7 +3863,7 @@
 
 // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK22: [[SIZE04:@.+]] = private {{.*}}constant [1 x i64] [i64 20]
-// CK22: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
+// CK22: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i64] [i64 51]
 
 // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK22: [[SIZE05:@.+]] = private {{.*}}constant [1 x i64] [i64 4]
@@ -3894,7 +3883,7 @@
 
 // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK22: [[SIZE09:@.+]] = private {{.*}}constant [1 x i64] [i64 20]
-// CK22: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
+// CK22: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 51]
 
 // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK22: [[SIZE10:@.+]] = private {{.*}}constant [1 x i64] [i64 4]
@@ -3914,7 +3903,7 @@
 
 // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK22: [[SIZE14:@.+]] = private {{.*}}constant [1 x i64] [i64 20]
-// CK22: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
+// CK22: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i64] [i64 51]
 
 int a;
 int c[100];
@@ -4010,11 +3999,10 @@
 
   // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-  // CK22-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**
+  // CK22-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32***
   // CK22-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
-  // CK22-DAG: store i32* [[RVAR0:%.+]], i32** [[CBP0]]
+  // CK22-DAG: store i32** @d, i32*** [[CBP0]]
   // CK22-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]]
-  // CK22-DAG: [[RVAR0]] = load i32*, i32** @d
   // CK22-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 2
   // CK22-DAG: [[RVAR00]] = load i32*, i32** @d
 
@@ -4093,11 +4081,10 @@
 
   // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-  // CK22-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
+  // CK22-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]***
   // CK22-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[ST]]**
-  // CK22-DAG: store [[ST]]* [[RVAR0:%.+]], [[ST]]** [[CBP0]]
+  // CK22-DAG: store [[ST]]** @sd, [[ST]]*** [[CBP0]]
   // CK22-DAG: store [[ST]]* [[SEC0:%.+]], [[ST]]** [[CP0]]
-  // CK22-DAG: [[RVAR0]] = load [[ST]]*, [[ST]]** @sd
   // CK22-DAG: [[SEC0]] = getelementptr {{.*}}[[ST]]* [[RVAR00:%.+]], i{{.+}} 2
   // CK22-DAG: [[RVAR00]] = load [[ST]]*, [[ST]]** @sd
 
@@ -4176,11 +4163,10 @@
 
   // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-  // CK22-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[STT]]**
+  // CK22-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[STT]]***
   // CK22-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[STT]]**
-  // CK22-DAG: store [[STT]]* [[RVAR0:%.+]], [[STT]]** [[CBP0]]
+  // CK22-DAG: store [[STT]]** @std, [[STT]]*** [[CBP0]]
   // CK22-DAG: store [[STT]]* [[SEC0:%.+]], [[STT]]** [[CP0]]
-  // CK22-DAG: [[RVAR0]] = load [[STT]]*, [[STT]]** @std
   // CK22-DAG: [[SEC0]] = getelementptr {{.*}}[[STT]]* [[RVAR00:%.+]], i{{.+}} 2
   // CK22-DAG: [[RVAR00]] = load [[STT]]*, [[STT]]** @std
 
@@ -4259,7 +4245,7 @@
 
 // CK23-LABEL: @.__omp_offloading_{{.*}}explicit_maps_inside_captured{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK23: [[SIZE05:@.+]] = private {{.*}}constant [1 x i64] [i64 16]
-// CK23: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
+// CK23: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i64] [i64 51]
 
 // CK23-LABEL: explicit_maps_inside_captured{{.*}}(
 int explicit_maps_inside_captured(int a){
@@ -4375,11 +4361,10 @@
 
     // CK23-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
     // CK23-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-    // CK23-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
+    // CK23-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float***
     // CK23-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
-    // CK23-DAG: store float* [[RVAR0:%.+]], float** [[CBP0]]
+    // CK23-DAG: store float** [[VAR0:%.+]], float*** [[CBP0]]
     // CK23-DAG: store float* [[SEC0:%.+]], float** [[CP0]]
-    // CK23-DAG: [[RVAR0]] = load float*, float** [[VAR0:%[^,]+]]
     // CK23-DAG: [[SEC0]] = getelementptr {{.*}}float* [[RVAR00:%.+]], i{{.+}} 2
     // CK23-DAG: [[RVAR00]] = load float*, float** [[VAR00:%[^,]+]]
     // CK23-DAG: [[VAR0]] = load float**, float*** [[CAP0:%[^,]+]]
@@ -5414,15 +5399,15 @@
 
 // CK27-LABEL: @.__omp_offloading_{{.*}}zero_size_section_and_private_maps{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK27: [[SIZE01:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer
-// CK27: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
+// CK27: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 51]
 
 // CK27-LABEL: @.__omp_offloading_{{.*}}zero_size_section_and_private_maps{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK27: [[SIZE02:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer
-// CK27: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
+// CK27: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i64] [i64 51]
 
 // CK27-LABEL: @.__omp_offloading_{{.*}}zero_size_section_and_private_maps{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK27: [[SIZE03:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer
-// CK27: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
+// CK27: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i64] [i64 51]
 
 // CK27-LABEL: @.__omp_offloading_{{.*}}zero_size_section_and_private_maps{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK27-LABEL: @.__omp_offloading_{{.*}}zero_size_section_and_private_maps{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
@@ -5470,11 +5455,10 @@
 
   // CK27-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK27-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-  // CK27-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**
+  // CK27-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32***
   // CK27-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
-  // CK27-DAG: store i32* [[RVAR0:%.+]], i32** [[CBP0]]
+  // CK27-DAG: store i32** [[VAR0:%.+]], i32*** [[CBP0]]
   // CK27-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]]
-  // CK27-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]]
   // CK27-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 0
   // CK27-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]]
 
@@ -5491,11 +5475,10 @@
 
   // CK27-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK27-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-  // CK27-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**
+  // CK27-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32***
   // CK27-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
-  // CK27-DAG: store i32* [[RVAR0:%.+]], i32** [[CBP0]]
+  // CK27-DAG: store i32** [[VAR0:%.+]], i32*** [[CBP0]]
   // CK27-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]]
-  // CK27-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]]
   // CK27-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 0
   // CK27-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]]
 
@@ -5512,11 +5495,10 @@
 
   // CK27-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK27-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-  // CK27-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**
+  // CK27-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32***
   // CK27-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
-  // CK27-DAG: store i32* [[RVAR0:%.+]], i32** [[CBP0]]
+  // CK27-DAG: store i32** [[VAR0:%.+]], i32*** [[CBP0]]
   // CK27-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]]
-  // CK27-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]]
   // CK27-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} %{{.+}}
   // CK27-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]]
 
@@ -5657,7 +5639,7 @@
 
 // CK28-LABEL: @.__omp_offloading_{{.*}}explicit_maps_pointer_references{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK28: [[SIZE01:@.+]] = private {{.*}}constant [1 x i64] [i64 400]
-// CK28: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
+// CK28: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 51]
 
 // CK28-LABEL: explicit_maps_pointer_references{{.*}}(
 void explicit_maps_pointer_references (int *p){
@@ -5690,11 +5672,10 @@
 
   // CK28-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK28-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-  // CK28-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**
+  // CK28-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32***
   // CK28-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
-  // CK28-DAG: store i32* [[VAR0:%.+]], i32** [[CBP0]]
+  // CK28-DAG: store i32** [[VAR00:%.+]], i32*** [[CBP0]]
   // CK28-DAG: store i32* [[VAR1:%.+]], i32** [[CP0]]
-  // CK28-DAG: [[VAR0]] = load i32*, i32** [[VAR00:%.+]],
   // CK28-DAG: [[VAR00]] = load i32**, i32*** [[VAR000:%.+]],
   // CK28-DAG: [[VAR1]] = getelementptr inbounds i32, i32* [[VAR11:%.+]], i{{64|32}} 2
   // CK28-DAG: [[VAR11]] = load i32*, i32** [[VAR111:%.+]],
@@ -6474,8 +6455,8 @@
 // SIMD-ONLY32-NOT: {{__kmpc|__tgt}}
 #ifdef CK32
 
-// CK32-DAG: [[MTYPE_TO:@.+]] = {{.+}}constant [1 x i64] [i64 33]
-// CK32-DAG: [[MTYPE_FROM:@.+]] = {{.+}}constant [1 x i64] [i64 34]
+// CK32-DAG: [[MTYPE_TO:@.+]] = {{.+}}constant [1 x i64] [i64 49]
+// CK32-DAG: [[MTYPE_FROM:@.+]] = {{.+}}constant [1 x i64] [i64 50]
 
 void array_shaping(float *f, int sa) {
 
@@ -6488,14 +6469,13 @@
   // CK32-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
   // CK32-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
 
-  // CK32-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to float**
+  // CK32-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to float***
   // CK32-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to float**
 
-  // CK32-DAG: store float* [[F1:%.+]], float** [[BPC0]],
+  // CK32-DAG: store float** [[F_ADDR:%.+]], float*** [[BPC0]],
   // CK32-DAG: store float* [[F2:%.+]], float** [[PC0]],
   // CK32-DAG: store i64 [[SIZE:%.+]], i64* [[S0]],
 
-  // CK32-DAG: [[F1]] = load float*, float** [[F_ADDR:%.+]],
   // CK32-DAG: [[F2]] = load float*, float** [[F_ADDR]],
   // CK32-64-DAG: [[SIZE]] = mul nuw i64 [[SZ1:%.+]], 4
   // CK32-64-DAG: [[SZ1]] = mul nuw i64 12, %{{.+}}
@@ -6514,14 +6494,13 @@
   // CK32-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
   // CK32-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
 
-  // CK32-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to float**
+  // CK32-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to float***
   // CK32-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to float**
 
-  // CK32-DAG: store float* [[F1:%.+]], float** [[BPC0]],
+  // CK32-DAG: store float** [[F_ADDR:%.+]], float*** [[BPC0]],
   // CK32-DAG: store float* [[F2:%.+]], float** [[PC0]],
   // CK32-DAG: store i64 [[SIZE:%.+]], i64* [[S0]],
 
-  // CK32-DAG: [[F1]] = load float*, float** [[F_ADDR:%.+]],
   // CK32-DAG: [[F2]] = load float*, float** [[F_ADDR]],
   // CK32-64-DAG: [[SIZE]] = mul nuw i64 [[SZ1:%.+]], 5
   // CK32-64-DAG: [[SZ1]] = mul nuw i64 4, %{{.+}}
diff --git a/clang/test/OpenMP/target_update_codegen.cpp b/clang/test/OpenMP/target_update_codegen.cpp
--- a/clang/test/OpenMP/target_update_codegen.cpp
+++ b/clang/test/OpenMP/target_update_codegen.cpp
@@ -310,18 +310,18 @@
 
 #ifdef CK5
 
-// CK5: [[SIZE00:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}, i{{64|32}} 4]
-// CK5: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 33, i64 17]
+// CK5: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
+// CK5: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 49]
 
 // CK5-LABEL: lvalue
 void lvalue(int *B, int l, int e) {
 
-  // CK5-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+  // CK5-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
   // CK5-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK5-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
-  // CK5-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
-  // CK5-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK5-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK5-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
   // CK5-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to i32***
   // CK5-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to i32**
   // CK5-DAG: store i32** [[B_ADDR:%.+]], i32*** [[BPC0]]
@@ -351,18 +351,18 @@
 
 #ifdef CK6
 
-// CK6: [[SIZE00:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}, i{{64|32}} 4]
-// CK6: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 33, i64 17]
+// CK6: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
+// CK6: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 49]
 
 // CK6-LABEL: lvalue
 void lvalue(int *B, int l, int e) {
 
-  // CK6-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+  // CK6-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
   // CK6-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK6-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
-  // CK6-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
-  // CK6-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK6-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK6-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
   // CK6-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to i32***
   // CK6-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to i32**
   // CK6-DAG: store i32** [[B_ADDR:%.+]], i32*** [[BPC0]]
@@ -397,18 +397,18 @@
 
 #ifdef CK7
 
-// CK7: [[SIZE00:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}, i{{64|32}} 4]
-// CK7: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 33, i64 17]
+// CK7: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
+// CK7: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 49]
 
 // CK7-LABEL: lvalue
 void lvalue(int *B, int l, int e) {
 
-  // CK7-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+  // CK7-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
   // CK7-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK7-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
-  // CK7-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
-  // CK7-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK7-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK7-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
   // CK7-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to i32***
   // CK7-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to i32**
   // CK7-DAG: store i32** [[B_ADDR:%.+]], i32*** [[BPC0]]
@@ -446,18 +446,18 @@
 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
 
 #ifdef CK8
-// CK8: [[SIZE00:@.+]] = {{.+}}constant [3 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}, i{{64|32}} {{8|4}}, i{{64|32}} 4]
-// CK8: [[MTYPE00:@.+]] = {{.+}}constant [3 x i64] [i64 33, i64 16, i64 17]
+// CK8: [[SIZE00:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}, i{{64|32}} 4]
+// CK8: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 49, i64 17]
 
 // CK8-LABEL: lvalue
 void lvalue(int **B, int l, int e) {
 
-  // CK8-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}], [3 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}, [3 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+  // CK8-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}], [2 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}, [2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
   // CK8-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK8-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
-  // CK8-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
-  // CK8-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
+  // CK8-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+  // CK8-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
   // CK8-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to i32***
   // CK8-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to i32**
   // CK8-DAG: store i32** [[ARRAY_IDX_1:%.+]], i32*** [[BPC0]]
@@ -466,13 +466,11 @@
   // CK8-DAG: [[ARRAY_IDX_1]] = getelementptr inbounds i32*, i32** [[ADD_PTR:%.+]], i{{.+}} 1
   // CK8-64-DAG: [[ADD_PTR]] = getelementptr inbounds i32*, i32** [[B_VAL:%.+]], i{{.+}} [[IDX_EXT:%.+]]
   // CK8-64-DAG: [[IDX_EXT]] = sext i32 [[L_VAL:%.+]] to i64
-  // CK8-DAG: [[ARRAY_IDX_4]] = getelementptr inbounds i32, i32* [[FIVE:%.+]], i{{.+}} 2
-  // CK8-64-DAG: [[ARRAY_IDX_3:%.+]] = getelementptr inbounds i32*, i32** [[ADD_PTR_2:%.+]], i{{.+}} 1
-  // CK8-64-DAG: [[ADD_PTR_2]] = getelementptr inbounds i32*, i32** %3, i{{.+}} [[IDX_EXT_1:%.+]]
-  // CK8-64-DAG: [[IDX_EXT_1]] = sext i32 [[L_VAL:%.+]] to i{{.+}}
+  // CK8-64-DAG: [[ARRAY_IDX_4]] = getelementptr inbounds i32, i32* [[FIVE:%.+]], i{{.+}} 2
   // CK8-32-DAG: [[ADD_PTR]] = getelementptr inbounds i32*, i32** [[B_VAL:%.+]], i{{.+}} [[L_VAL:%.+]]
-  // CK8-32-DAG: [[ARRAY_IDX_4:%.+]] = getelementptr inbounds i32*, i32** [[ADD_PTR_2:%.+]], i{{.+}} 1
-  // CK8-32-DAG: [[ADD_PTR_2]] = getelementptr inbounds i32*, i32** %3, i{{.+}} [[L_VAL:%.+]]
+  // CK8-32-DAG: [[ARRAY_IDX_4:%.+]] = getelementptr inbounds i32, i32* [[ADD_PTR_2:%.+]], i{{.+}} 2
+  // CK8-32-DAG: [[ADD_PTR_2]] = load i32*, i32** [[ARRAY_IDX_2:%.+]],
+  // CK8-32-DAG: [[ARRAY_IDX_2]]  = getelementptr inbounds i32*, i32** %{{.+}}, i32 1
   #pragma omp target update to((B+l)[1][2])
   (B+l)[1][2] += e;
   #pragma omp target update from((B+l)[1][2])
@@ -501,19 +499,19 @@
   double *p;
 };
 
-// CK9: [[MTYPE00:@.+]] = {{.+}}constant [3 x i64] [i64 32, i64 281474976710657, i64 281474976710673]
+// CK9: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673]
 
 // CK9-LABEL: lvalue
 void lvalue(struct S *s, int l, int e) {
 
-  // CK9-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}, [3 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+  // CK9-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}, [2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
   // CK9-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK9-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
   // CK9-DAG: [[GSIZE]] = getelementptr inbounds {{.+}}[[SIZE:%[^,]+]]
   //
-  // CK9-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
-  // CK9-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
-  // CK9-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 2
+  // CK9-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+  // CK9-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK9-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1
   // CK9-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to double***
   // CK9-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to double**
   // CK9-DAG: store double** [[P:%.+]], double*** [[BPC0]]
@@ -551,19 +549,19 @@
   double *p;
 };
 
-// CK10: [[MTYPE00:@.+]] = {{.+}}constant [3 x i64] [i64 32, i64 281474976710657, i64 281474976710673]
+// CK10: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673]
 
 // CK10-LABEL: lvalue
 void lvalue(struct S *s, int l, int e) {
 
-  // CK10-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}, [3 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+  // CK10-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}, [2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
   // CK10-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK10-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
   // CK10-DAG: [[GSIZE]] = getelementptr inbounds {{.+}}[[SIZE:%[^,]+]]
   //
-  // CK10-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
-  // CK10-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
-  // CK10-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 2
+  // CK10-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+  // CK10-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK10-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1
   // CK10-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to double***
   // CK10-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to double**
   // CK10-DAG: store double** [[P_VAL:%.+]], double*** [[BPC0]]
@@ -601,19 +599,19 @@
 struct S {
   double *p;
 };
-// CK11: [[MTYPE00:@.+]] = {{.+}}constant [3 x i64] [i64 32, i64 281474976710657, i64 281474976710673]
+// CK11: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673]
 
 // CK11-LABEL: lvalue
 void lvalue(struct S *s, int l, int e) {
 
-  // CK11-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}, [3 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+  // CK11-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}, [2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
   // CK11-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK11-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
   // CK11-DAG: [[GSIZE]] = getelementptr inbounds {{.+}}[[SIZE:%[^,]+]]
   //
-  // CK11-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
-  // CK11-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
-  // CK11-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 2
+  // CK11-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+  // CK11-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK11-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1
   // CK11-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to double***
   // CK11-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to double**
   // CK11-DAG: store double** [[P:%.+]], double*** [[BPC0]]
@@ -653,41 +651,41 @@
   double *p;
   struct S *sp;
 };
-// CK12: [[MTYPE00:@.+]] = {{.+}}constant [4 x i64] [i64 32, i64 281474976710657, i64 281474976710672, i64 17]
+// CK12: [[MTYPE00:@.+]] = {{.+}}constant [3 x i64] [i64 32, i64 281474976710672, i64 17]
 
 // CK12-LABEL: lvalue
 void lvalue(struct S *s, int l, int e) {
 
-  // CK12-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 4, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[4 x i{{.+}}, [4 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+  // CK12-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}, [3 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
   // CK12-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK12-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
   // CK12-DAG: [[GSIZE]] = getelementptr inbounds {{.+}}[[SIZE:%[^,]+]]
   //
-  // CK12-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3
-  // CK12-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3
-  // CK12-DAG: [[SIZE2:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 3
+  // CK12-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
+  // CK12-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
+  // CK12-DAG: [[SIZE2:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 2
   // CK12-DAG: [[BPC2:%.+]] = bitcast i8** [[BP2]] to double***
   // CK12-DAG: [[PC2:%.+]] = bitcast i8** [[P2]] to double**
   // CK12-DAG: store double** [[P_VAL:%.+]], double*** [[BPC2]]
   // CK12-DAG: store double* [[SIX:%.+]], double** [[PC2]]
   // CK12-DAG: store i{{.+}} 8, i{{.+}}* [[SIZE2]]
-  // CK12-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
-  // CK12-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
-  // CK12-DAG: [[SIZE1:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 2
+  // CK12-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+  // CK12-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK12-DAG: [[SIZE1:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1
   // CK12-DAG: [[BPC1:%.+]] = bitcast i8** [[BP1]] to [[STRUCT_S:%.+]]***
   // CK12-DAG: [[PC1:%.+]] = bitcast i8** [[P1]] to double***
   // CK12-DAG: store [[STRUCT_S]]** [[SP:%.+]], [[STRUCT_S]]*** [[BPC1]]
   // CK12-DAG: store double** [[P_VAL:%.+]], double*** [[PC1]]
   // CK12-DAG: store i{{.+}} {{4|8}}, i{{.+}}* [[SIZE1]]
-  // CK12-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
-  // CK12-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
-  // CK12-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1
+  // CK12-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK12-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK12-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 0
   // CK12-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to [[STRUCT_S:%.+]]**
   // CK12-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to [[STRUCT_S]]***
-  // CK12-DAG: store [[STRUCT_S]]** [[S:%.+]], [[STRUCT_S]]*** [[S_VAL:%.+]]
-  // CK12-DAG: store i{{.+}} {{.+}}, i{{.+}}* [[SIZE0]]
+  // CK12-DAG: store [[STRUCT_S]]* [[ZERO:%.+]], [[STRUCT_S]]** [[BPC0]]
   // CK12-DAG: [[SP]] = getelementptr inbounds [[STRUCT_S]], [[STRUCT_S]]* [[ONE:%.+]], i32 0, i32 1
-  // CK12-DAG: [[ONE]] = load %struct.S*, %struct.S** [[S]],
+  // CK12-DAG: [[ONE]] = load %struct.S*, %struct.S** [[S:%.+]],
+  // CK12-DAG: [[ZERO]] = load %struct.S*, %struct.S** [[S]],
   #pragma omp target update to(*(s->sp->p))
     *(s->sp->p) = e;
   #pragma omp target update from(*(s->sp->p))
@@ -711,18 +709,18 @@
 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
 #ifdef CK13
 
-// CK13: [[SIZE00:@.+]] = {{.+}}constant [2 x i64] [i64 {{8|4}}, i64 4]
-// CK13: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 33, i64 17]
+// CK13: [[SIZE00:@.+]] = {{.+}}constant [1 x i64] [i64 4]
+// CK13: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 49]
 
 // CK13-LABEL: lvalue
 void lvalue(int **BB, int a, int b) {
 
-  // CK13-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+  // CK13-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
   // CK13-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK13-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
-  // CK13-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
-  // CK13-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK13-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK13-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
   // CK13-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to i32****
   // CK13-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to i32**
   // CK13-DAG: store i32*** [[BB_ADDR:%.+]], i32**** [[BPC0]]
@@ -831,7 +829,7 @@
 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
 #ifdef CK15
 
-// CK15: [[MTYPE00:@.+]] = {{.+}}constant [3 x i64] [i64 32, i64 281474976710657, i64 281474976710673]
+// CK15: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673]
 
 struct SSA {
   double *p;
@@ -842,14 +840,14 @@
 //CK-15-LABEL: lvalue_member
 void lvalue_member(SSA *sap) {
 
-  // CK15-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+  // CK15-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
   // CK15-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK15-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
   // CK15-DAG: [[GSIZE]] = getelementptr inbounds {{.+}}[[SIZE:%[^,]+]]
 
-  // CK15-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
-  // CK15-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
-  // CK15-DAG: [[SIZE2:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 2
+  // CK15-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+  // CK15-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK15-DAG: [[SIZE2:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1
   // CK15-DAG: [[BPC2:%.+]] = bitcast i8** [[BP2]] to double***
   // CK15-DAG: [[PC2:%.+]] = bitcast i8** [[P2]] to double**
   // CK15-DAG: store double** [[P_VAL:%.+]], double*** [[BPC2]]
@@ -858,20 +856,11 @@
   // CK15-DAG: [[ADD_PTR]] = getelementptr inbounds double, double* [[THREE:%.+]], i{{.+}} 3
   // CK15-DAG: [[THREE]] = load double*, double** [[P_VAL_1:%.+]]
   // CK15-DAG: [[P_VAL]] = getelementptr inbounds [[SSA:%.+]], [[SSA:%.+]]* [[THIS:%.+]], i32 0, i32 0
-  // CK15-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
-  // CK15-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
-  // CK15-DAG: [[SIZE1:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1
-  // CK15-DAG: [[BPC1:%.+]] = bitcast i8** [[BP1]] to [[SSA]]**
-  // CK15-DAG: [[PC1:%.+]] = bitcast i8** [[P1]] to [[SSA]]***
-  // CK15-DAG: store [[SSA]]* [[SAP_VAL:%.+]], [[SSA]]** [[BPC1]],
-  // CK15-DAG: store [[SSA]]** [[SAP_ADDR:%.+]], [[SSA]]*** [[PC1]]
-  // CK15-DAG: store i{{.+}} {{8|4}}, i{{.+}}* [[SIZE1]]
-  // CK15-DAG: [[SAP_VAL]] = load [[SSA]]*, [[SSA]]** [[SAP_ADDR]],
   // CK15-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // CK15-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
   // CK15-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 0
-  // CK15-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to [[SSA]]***
-  // CK15-DAG: store [[SSA]]** [[SAP_ADDR]], [[SSA]]*** [[BPC0]],
+  // CK15-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to [[SSA]]**
+  // CK15-DAG: store [[SSA]]* [[SAP_ADDR:%.+]], [[SSA]]** [[BPC0]],
   // CK15-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to double***
   // CK15-DAG: store double** [[P_VAL]], double*** [[PC0]],
   // CK15-DAG: store i{{.+}} [[COMPUTE_SIZE:%.+]], i{{.+}}* [[SIZE0]]
@@ -904,18 +893,18 @@
 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
 #ifdef CK16
 
-// CK16: [[SIZE00:@.+]] = {{.+}}constant [2 x i64] [i64 {{8|4}}, i64 4]
-// CK16: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 33, i64 17]
+// CK16: [[SIZE00:@.+]] = {{.+}}constant [1 x i64] [i64 4]
+// CK16: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 49]
 
 //CK16-LABEL: lvalue_find_base
 void lvalue_find_base(float *f, int *i) {
 
-  // CK16-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+  // CK16-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
   // CK16-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK16-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
-  // CK16-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
-  // CK16-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK16-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK16-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
   // CK16-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to float***
   // CK16-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to float**
   // CK16-DAG: store float** [[F_ADDR:%.+]], float*** [[BPC0]]
@@ -948,8 +937,8 @@
 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
 #ifdef CK17
 
-// CK17: [[SIZE00:@.+]] = {{.+}}constant [2 x i64] [i64 {{4|8}}, i64 4]
-// CK17: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 33, i64 17]
+// CK17: [[SIZE00:@.+]] = {{.+}}constant [1 x i64] [i64 4]
+// CK17: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 49]
 
 struct SSA {
   int i;
@@ -959,12 +948,12 @@
 //CK17-LABEL: lvalue_find_base
 void lvalue_find_base(float **f, SSA *sa) {
 
-  // CK17-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+  // CK17-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
   // CK17-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK17-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
-  // CK17-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
-  // CK17-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK17-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK17-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
   // CK17-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to float****
   // CK17-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to float**
   // CK17-DAG: store float*** [[F_ADDR:%.+]], float**** [[BPC0]],
@@ -1005,37 +994,25 @@
 // SIMD-ONLY18-NOT: {{__kmpc|__tgt}}
 #ifdef CK18
 
-// CK18-DAG: [[MTYPE_TO:@.+]] = {{.+}}constant [2 x i64] [i64 33, i64 16]
-// CK18-DAG: [[MTYPE_FROM:@.+]] = {{.+}}constant [2 x i64] [i64 34, i64 16]
+// CK18-DAG: [[MTYPE_TO:@.+]] = {{.+}}constant [1 x i64] [i64 49]
+// CK18-DAG: [[MTYPE_FROM:@.+]] = {{.+}}constant [1 x i64] [i64 50]
 
 //CK18-LABEL: array_shaping
 void array_shaping(float *f, int sa) {
 
-  // CK18-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE_TO]]{{.+}}, i8** null)
+  // CK18-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE_TO]]{{.+}}, i8** null)
   // CK18-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK18-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
   // CK18-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 
-  // CK18-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
-  // CK18-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-  // CK18-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
-
-  // CK18-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to float**
-  // CK18-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to float***
-
-  // CK18-DAG: store float* [[F1:%.+]], float** [[BPC0]],
-  // CK18-DAG: store float** [[F_ADDR:%.+]], float*** [[PC0]],
-  // CK18-DAG: store i64 {{8|4}}, i64* [[S0]],
-  // CK18-DAG: [[F1]] = load float*, float** [[F_ADDR]],
-
-  // CK18-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
-  // CK18-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
-  // CK18-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
+  // CK18-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK18-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK18-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
 
   // CK18-DAG: [[BPC1:%.+]] = bitcast i8** [[BP1]] to float***
   // CK18-DAG: [[PC1:%.+]] = bitcast i8** [[P1]] to float**
 
-  // CK18-DAG: store float** [[F_ADDR]], float*** [[BPC1]],
+  // CK18-DAG: store float** [[F_ADDR:%.+]], float*** [[BPC1]],
   // CK18-DAG: store float* [[F2:%.+]], float** [[PC1]],
   // CK18-DAG: store i64 [[SIZE:%.+]], i64* [[S1]],
   // CK18-DAG: [[F2]] = load float*, float** [[F_ADDR]],
@@ -1047,31 +1024,19 @@
   // CK18-32-DAG: [[SZ2]] = mul nuw i32 12, %{{.+}}
   #pragma omp target update to(([3][sa][4])f)
   sa = 1;
-  // CK18-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE_FROM]]{{.+}}, i8** null)
+  // CK18-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE_FROM]]{{.+}}, i8** null)
   // CK18-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK18-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
   // CK18-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 
-  // CK18-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
-  // CK18-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-  // CK18-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
-
-  // CK18-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to float**
-  // CK18-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to float***
-
-  // CK18-DAG: store float* [[F1:%.+]], float** [[BPC0]],
-  // CK18-DAG: store float** [[F_ADDR:%.+]], float*** [[PC0]],
-  // CK18-DAG: store i64 {{8|4}}, i64* [[S0]],
-  // CK18-DAG: [[F1]] = load float*, float** [[F_ADDR]],
-
-  // CK18-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
-  // CK18-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
-  // CK18-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
+  // CK18-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK18-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK18-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
 
   // CK18-DAG: [[BPC1:%.+]] = bitcast i8** [[BP1]] to float***
   // CK18-DAG: [[PC1:%.+]] = bitcast i8** [[P1]] to float**
 
-  // CK18-DAG: store float** [[F_ADDR]], float*** [[BPC1]],
+  // CK18-DAG: store float** [[F_ADDR:%.+]], float*** [[BPC1]],
   // CK18-DAG: store float* [[F2:%.+]], float** [[PC1]],
   // CK18-DAG: store i64 [[SIZE:%.+]], i64* [[S1]],
   // CK18-DAG: [[F2]] = load float*, float** [[F_ADDR]],
diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp
--- a/openmp/libomptarget/src/omptarget.cpp
+++ b/openmp/libomptarget/src/omptarget.cpp
@@ -872,14 +872,9 @@
           return OFFLOAD_FAIL;
         }
       }
-    } else if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ) {
-      TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBase, sizeof(void *), IsLast,
-          false, IsHostPtr);
-      TgtBaseOffset = 0; // no offset for ptrs.
-      DP("Obtained target argument " DPxMOD " from host pointer " DPxMOD " to "
-         "object " DPxMOD "\n", DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBase),
-         DPxPTR(HstPtrBase));
     } else {
+      if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)
+        HstPtrBase = *reinterpret_cast<void**>(HstPtrBase);
       TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, arg_sizes[i], IsLast,
           false, IsHostPtr);
       TgtBaseOffset = (intptr_t)HstPtrBase - (intptr_t)HstPtrBegin;