[OPENMP]Fix PR46824: Global declare target pointer cannot be accessed in target region.

[ABataev]

Summary

Need to map the base pointer for all directives, not only target
data-based ones.
The base pointer is mapped for array sections, array subscript, array
shaping and other array-like constructs with the base pointer. Also,
codegen for use_device_ptr clause was modified to correctly handle
mapping combination of array like constructs + use_device_ptr clause.
The data for use_device_ptr clause is emitted as the last records in the
data mapping array.

[ye-luo]

Please check the reproducer in https://bugs.llvm.org/show_bug.cgi?id=46868 with LIBOMPTARGET_DEBUG=1.
The reference counting on the base pointer variable has side effects. It was not cleaned up when these variables leave its scope.

[ABataev]

Fixed ref count for base pointer + extra test.

[ye-luo]

This patch
GPU activities: 96.99% 350.05ms 10 35.005ms 1.5680us 350.00ms [CUDA memcpy HtoD]
before the July21 change
GPU activities: 95.33% 20.317ms 4 5.0793ms 1.6000us 20.305ms [CUDA memcpy HtoD]
Still more transfer than it should.

[ye-luo]

In D84767#2180280, @ye-luo wrote:

This patch
GPU activities: 96.99% 350.05ms 10 35.005ms 1.5680us 350.00ms [CUDA memcpy HtoD]
before the July21 change
GPU activities: 95.33% 20.317ms 4 5.0793ms 1.6000us 20.305ms [CUDA memcpy HtoD]
Still more transfer than it should.

@ABataev could you have a look? My July 21 compiler was built before "[OPENMP]Fix PR46012: declare target pointer cannot be accessed in target region." gets in.

[ABataev]

In D84767#2180433, @ye-luo wrote:

In D84767#2180280, @ye-luo wrote:

This patch
GPU activities: 96.99% 350.05ms 10 35.005ms 1.5680us 350.00ms [CUDA memcpy HtoD]
before the July21 change
GPU activities: 95.33% 20.317ms 4 5.0793ms 1.6000us 20.305ms [CUDA memcpy HtoD]
Still more transfer than it should.

@ABataev could you have a look? My July 21 compiler was built before "[OPENMP]Fix PR46012: declare target pointer cannot be accessed in target region." gets in.

Are you talking about the number of calls value? The total number of calls will increase after the patch anyway, PTR_AND_OBJ adds 1 extra mem transfer for transferring translated pointer address.

[ABataev]

Rebase + apply to global pointers only.

[grokos]

clang/test/OpenMP/target_data_codegen.cpp

659–660

MEMBER_OF_7=0x9000000000000 --> MEMBER_OF_7=0x7000000000000

[ABataev]

Fixed comment in test

[ye-luo]

LGTM. My applications run as expected now. PR46824, PR46012, PR46868 all work fine.

clang/lib/CodeGen/CGOpenMPRuntime.cpp

void generateInfoForComponentList(
     //
     // map(i[1:23])
     // &i(=&i[0]), &i[1], 23*sizeof(int), TARGET_PARAM | TO | FROM
     //
     // map(p)
     // &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
     //
     // map(s.i)
     // &s, &(s.i), sizeof(int), TARGET_PARAM | TO | FROM
     //
     // map(s.s.f)
void generateInfoForComponentList(
     // Scan the components from the base to the complete expression.
     auto CI = Components.rbegin();
     auto CE = Components.rend();
     auto I = CI;
 
     // 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);
     const auto *OASE = dyn_cast<OMPArraySectionExpr>(AssocExpr);
     const auto *OAShE = dyn_cast<OMPArrayShapingExpr>(AssocExpr);
 
     if (isa<MemberExpr>(AssocExpr)) {
       // The base is the 'this' pointer. The content of the pointer is going
if (isa<MemberExpr>(AssocExpr)) {
       }
 
       // If the variable is a pointer and is being dereferenced (i.e. is not
       // the last component), the base has to be the pointer itself, not its
       // reference. References are ignored for mapping purposes.
       QualType Ty =
           I->getAssociatedDeclaration()->getType().getNonReferenceType();
       if (Ty->isAnyPointerType() && std::next(I) != CE) {
+        // 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>());
-
-        // 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;
+        else
+          FirstPointerInComplexData = true;
       }
     }
 
     // Track whether a component of the list should be marked as MEMBER_OF some
     // combined entry (for partial structs). Only the first PTR_AND_OBJ entry
     // in a component list should be marked as MEMBER_OF, all subsequent entries
     // do not belong to the base struct. E.g.
     // struct S2 s;
void generateInfoForComponentList(
     const MemberExpr *EncounteredME = nullptr;
 
     for (; I != CE; ++I) {
       // If the current component is member of a struct (parent struct) mark it.
       if (!EncounteredME) {
         EncounteredME = dyn_cast<MemberExpr>(I->getAssociatedExpression());
         // If we encounter a PTR_AND_OBJ entry from now on it should be marked
         // as MEMBER_OF the parent struct.
-        if (EncounteredME)
+        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);
 
       // We need to generate the addresses and sizes if this is the last
       // component, if the component is a pointer or if it is an array section
       // whose length can't be proved to be one. If this is a pointer, it
       // becomes the base address for the following components.
for (; I != CE; ++I) {
           // If Mapper is valid, the last component inherits the mapper.
           bool HasMapper = Mapper && Next == CE;
           CombinedInfo.Mappers.push_back(HasMapper ? Mapper : nullptr);
 
           // We need to add a pointer flag for each map that comes from the
           // same expression except for the first one. We also need to signal
           // this map is the first one that relates with the current capture
           // (there is a set of entries for each capture).
-          OpenMPOffloadMappingFlags Flags = getMapTypeBits(
-              MapType, MapModifiers, IsImplicit,
-              !IsExpressionFirstInfo || 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,
             // then we reset the TO/FROM/ALWAYS/DELETE/CLOSE flags.
             if (IsPointer)
               Flags &= ~(OMP_MAP_TO | OMP_MAP_FROM | OMP_MAP_ALWAYS |
                          OMP_MAP_DELETE | OMP_MAP_CLOSE);
 
for (; I != CE; ++I) {
           break;
 
         // The pointer becomes the base for the next element.
         if (Next != CE)
           BP = LB;
 
         IsExpressionFirstInfo = false;
         IsCaptureFirstInfo = false;
+        FirstPointerInComplexData = false;
       }
     }
   }
 
   /// Return the adjusted map modifiers if the declaration a capture refers to
   /// appears in a first-private clause. This is expected to be used only with
   /// directives that start with 'target'.
   MappableExprsHandler::OpenMPOffloadMappingFlags
void generateAllInfo(
     // Look at the use_device_ptr clause information and mark the existing map
     // entries as such. If there is no map information for an entry in the
     // use_device_ptr list, we create one with map type 'alloc' and zero size
     // section. It is the user fault if that was not mapped before. If there is
     // no map information and the pointer is a struct member, then we defer the
     // 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>()) {
       for (const auto L : C->component_lists()) {
         OMPClauseMappableExprCommon::MappableExprComponentListRef Components =
             std::get<1>(L);
         assert(!Components.empty() &&
                "Not expecting empty list of components!");
         const ValueDecl *VD = Components.back().getAssociatedDeclaration();
         VD = cast<ValueDecl>(VD->getCanonicalDecl());
         const Expr *IE = Components.back().getAssociatedExpression();
         // If the first component is a member expression, we have to look into
         // 'this', which maps to null in the map of map information. Otherwise
         // look directly for the information.
         auto It = Info.find(isa<MemberExpr>(IE) ? nullptr : VD);
 
         // 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) {
           // 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;
           }
         }
 
         // We didn't find any match in our map information - generate a zero
         // size array section - if the pointer is a struct member we defer this
         // action until the whole struct has been processed.
         if (isa<MemberExpr>(IE)) {
           // Insert the pointer into Info to be processed by
           // generateInfoForComponentList. Because it is a member pointer
           // without a pointee, no entry will be generated for it, therefore
           // we need to generate one after the whole struct has been processed.
           // Nonetheless, generateInfoForComponentList must be called to take
           // the pointer into account for the calculation of the range of the
           // partial struct.
           InfoGen(nullptr, Components, OMPC_MAP_unknown, llvm::None,
                   /*ReturnDevicePointer=*/false, C->isImplicit(), nullptr);
           DeferredInfo[nullptr].emplace_back(IE, VD, /*ForDeviceAddr=*/false);
         } 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 |
+          UseDevicePtrCombinedInfo.Types.push_back(OMP_MAP_RETURN_PARAM |
                                                    OMP_MAP_TARGET_PARAM);
-          CombinedInfo.Mappers.push_back(nullptr);
+          UseDevicePtrCombinedInfo.Mappers.push_back(nullptr);
         }
       }
     }
 
     // Look at the use_device_addr clause information and mark the existing map
     // entries as such. If there is no map information for an entry in the
     // use_device_addr list, we create one with map type 'alloc' and zero size
     // section. It is the user fault if that was not mapped before. If there is
for (const auto &M : Info) {
       // If there is an entry in PartialStruct it means we have a struct with
       // individual members mapped. Emit an extra combined entry.
       if (PartialStruct.Base.isValid())
         emitCombinedEntry(CombinedInfo, CurInfo.Types, PartialStruct);
 
       // 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
   /// mappers for the extracted map clauses of user-defined mapper (all included
   /// in \a CombinedInfo).
   void generateAllInfoForMapper(MapCombinedInfoTy &CombinedInfo) const {
     assert(CurDir.is<const OMPDeclareMapperDecl *>() &&
            "Expect a declare mapper directive");

clang/test/OpenMP/target_data_codegen.cpp

 // expected-no-diagnostics
 #ifndef HEADER
 #define HEADER
 
 ///==========================================================================///
 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-64
 // RUN: %clang_cc1 -DCK1 -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 CK1 --check-prefix CK1-64
 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK1 --check-prefix CK1-32

 };
 struct S2 {
   S1 s;
   struct S2 *ps;
 };
 
 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);
   }
 }
 #endif
 ///==========================================================================///
 // RUN: %clang_cc1 -DCK6 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK6 --check-prefix CK6-64

 
 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.
 
-  // 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
   //
   // PRESENT=0x1000 | TARGET_PARAM=0x20 | FROM=0x2 | TO=0x1 = 0x1023
   //
   // CK8-SAME: {{^}} i64 [[#0x1023]],
 
   // 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)
   {
     ++(arg);
   }
 }
 #endif

clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp

 // expected-no-diagnostics
 #ifndef HEADER
 #define HEADER
 
 ///==========================================================================///
 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-64
 // RUN: %clang_cc1 -DCK1 -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 CK1 --check-prefix CK1-64
 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK1 --check-prefix CK1-32
 // RUN: %clang_cc1 -DCK1 -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 CK1 --check-prefix CK1-32
 
 // RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // RUN: %clang_cc1 -DCK1 -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 --check-prefix SIMD-ONLY0 %s
 // RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // RUN: %clang_cc1 -DCK1 -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 --check-prefix SIMD-ONLY0 %s
 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
 #ifdef CK1
 
 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>
 void foo(float *&lr, T *&tr) {
   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
   #pragma omp target data map(g[:10]) use_device_ptr(g)
   {
     ++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]],
   // CK1:     store float* [[VAL]], float** [[PVT:%.+]],
   // CK1:     [[TT1:%.+]] = load float*, float** [[PVT]],
   // CK1:     getelementptr inbounds float, float* [[TT1]], i32 1
   #pragma omp target data map(l[:10]) use_device_ptr(l)
   {
void foo(float *&lr, T *&tr) {
   {
     ++l;
   }
   // CK1-NOT: call void @__tgt_target
   // CK1:     [[TTT:%.+]] = load float*, float** [[DECL]],
   // 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]],
   // CK1:     store float* [[VAL]], float** [[PVT:%.+]],
   // CK1:     [[TT1:%.+]] = load float*, float** [[PVT]],
   // CK1:     getelementptr inbounds float, float* [[TT1]], i32 1
   #pragma omp target data map(l[:10]) use_device_ptr(l) if(1)
   {
     ++l;
   }
   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE03]]
   // CK1:     [[TTT:%.+]] = load float*, float** [[DECL]],
   // CK1:     getelementptr inbounds float, float* [[TTT]], i32 1
   ++l;
 
   // CK1:     [[CMP:%.+]] = icmp ne float* %{{.+}}, null
   // 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]],
   // CK1:     store float* [[VAL]], float** [[PVT:%.+]],
   // CK1:     [[TT1:%.+]] = load float*, float** [[PVT]],
   // CK1:     getelementptr inbounds float, float* [[TT1]], i32 1
   // CK1:     br label %[[BEND:.+]]
 
void foo(float *&lr, T *&tr) {
   // CK1:     br label %[[BEND]]
 
   // CK1:     [[BEND]]:
   // CK1:     [[TTT:%.+]] = load float*, float** [[DECL]],
   // CK1:     getelementptr inbounds float, float* [[TTT]], i32 1
   ++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:%.+]],
   // CK1-NOT: store float** [[PVTV]], float*** [[DECL]],
   // CK1:     store float** [[PVTV]], float*** [[PVT:%.+]],
   // CK1:     [[TT1:%.+]] = load float**, float*** [[PVT]],
   // CK1:     [[TT2:%.+]] = load float*, float** [[TT1]],
   // CK1:     getelementptr inbounds float, float* [[TT2]], i32 1
   #pragma omp target data map(lr[:10]) use_device_ptr(lr)
   {
     ++lr;
   }
   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE05]]
   // CK1:     [[TTT:%.+]] = load float**, float*** [[DECL]],
   // CK1:     [[TTTT:%.+]] = load float*, float** [[TTT]],
   // 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]],
   // CK1:     store i32* [[VAL]], i32** [[PVT:%.+]],
   // CK1:     [[TT1:%.+]] = load i32*, i32** [[PVT]],
   // CK1:     getelementptr inbounds i32, i32* [[TT1]], i32 1
   #pragma omp target data map(t[:10]) use_device_ptr(t)
   {
     ++t;
   }
   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE06]]
   // CK1:     [[TTT:%.+]] = load i32*, i32** [[DECL]],
   // 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:%.+]],
   // CK1-NOT: store i32** [[PVTV]], i32*** [[DECL]],
   // CK1:     store i32** [[PVTV]], i32*** [[PVT:%.+]],
   // CK1:     [[TT1:%.+]] = load i32**, i32*** [[PVT]],
   // CK1:     [[TT2:%.+]] = load i32*, i32** [[TT1]],
   // CK1:     getelementptr inbounds i32, i32* [[TT2]], i32 1
   #pragma omp target data map(tr[:10]) use_device_ptr(tr)
   {
     ++tr;
   }
   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE07]]
   // CK1:     [[TTT:%.+]] = load i32**, i32*** [[DECL]],
   // CK1:     [[TTTT:%.+]] = load i32*, i32** [[TTT]],
   // 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]],
   // CK1:     store float* [[VAL]], float** [[PVT:%.+]],
   // CK1:     [[TT1:%.+]] = load float*, float** [[PVT]],
   // CK1:     getelementptr inbounds float, float* [[TT1]], i32 1
   #pragma omp target data map(l[:10], t[:10]) use_device_ptr(l)
   {
     ++l; ++t;
   }
   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE08]]
   // CK1:     [[TTT:%.+]] = load float*, float** [[DECL]],
   // CK1:     getelementptr inbounds float, float* [[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{{.+}}[[MTYPE09]]
   // CK1:     [[_VAL:%.+]] = load float*, float** [[_CBP]],
   // CK1:     store float* [[_VAL]], float** [[_PVT:%.+]],
   // CK1:     [[VAL:%.+]] = load i32*, i32** [[CBP]],
   // CK1:     store i32* [[VAL]], i32** [[PVT:%.+]],
   // CK1:     [[_TT1:%.+]] = load float*, float** [[_PVT]],
   // CK1:     getelementptr inbounds float, float* [[_TT1]], i32 1
   // CK1:     [[TT1:%.+]] = load i32*, i32** [[PVT]],
   // CK1:     getelementptr inbounds i32, i32* [[TT1]], i32 1
   #pragma omp target data map(l[:10], t[:10]) use_device_ptr(l) use_device_ptr(t)
   {
     ++l; ++t;
   }
   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE09]]
   // CK1:     [[_TTT:%.+]] = load float*, float** {{%.+}},
   // CK1:     getelementptr inbounds float, float* [[_TTT]], i32 1
   // CK1:     [[TTT:%.+]] = load i32*, i32** {{%.+}},
   // 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]]
   // CK1:     [[_VAL:%.+]] = load float*, float** [[_CBP]],
   // CK1:     store float* [[_VAL]], float** [[_PVT:%.+]],
   // CK1:     [[VAL:%.+]] = load i32*, i32** [[CBP]],
   // CK1:     store i32* [[VAL]], i32** [[PVT:%.+]],
   // CK1:     [[_TT1:%.+]] = load float*, float** [[_PVT]],
   // CK1:     getelementptr inbounds float, float* [[_TT1]], i32 1
   // CK1:     [[TT1:%.+]] = load i32*, i32** [[PVT]],
   // CK1:     getelementptr inbounds i32, i32* [[TT1]], i32 1
   #pragma omp target data map(l[:10], t[:10]) use_device_ptr(l,t)
   {
     ++l; ++t;
   }
   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE10]]
   // CK1:     [[_TTT:%.+]] = load float*, float** {{%.+}},
   // CK1:     getelementptr inbounds float, float* [[_TTT]], i32 1
   // CK1:     [[TTT:%.+]] = load i32*, i32** {{%.+}},
   // CK1:     getelementptr inbounds i32, i32* [[TTT]], i32 1
   ++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]]
   // CK1:     [[VAL:%.+]] = load i32*, i32** [[CBP]],
   // CK1-NOT: store i32* [[VAL]], i32** [[DECL]],
   // CK1:     store i32* [[VAL]], i32** [[PVT:%.+]],
   // CK1:     [[TT1:%.+]] = load i32*, i32** [[PVT]],
   // CK1:     getelementptr inbounds i32, i32* [[TT1]], i32 1
   #pragma omp target data map(l[:10]) use_device_ptr(t)
   {
     ++l; ++t;
   }
   // CK1:     call void @__tgt_target_data_end{{.+}}[[MTYPE11]]
   // CK1:     [[TTT:%.+]] = load i32*, i32** [[DECL]],
   // CK1:     getelementptr inbounds i32, i32* [[TTT]], i32 1
   ++l; ++t;
 
   // 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]]
   // CK1:     [[VAL:%.+]] = load i32*, i32** [[CBP]],
   // CK1:     store i32* [[VAL]], i32** [[PVTV:%.+]],
   // CK1-NOT: store i32** [[PVTV]], i32*** [[DECL]],
   // CK1:     store i32** [[PVTV]], i32*** [[PVT:%.+]],
   // CK1:     [[TT1:%.+]] = load i32**, i32*** [[PVT]],

 // RUN: %clang_cc1 -DCK2 -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 --check-prefix SIMD-ONLY1 %s
 // SIMD-ONLY1-NOT: {{__kmpc|__tgt}}
 #ifdef CK2
 
 // 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>
 struct ST {
   T *a;
   double *&b;
   ST(double *&b) : a(0), b(b) {}
 
void foo(double *&arg) {
     }
     // CK2:     call void @__tgt_target_data_end{{.+}}[[MTYPE01]]
     // CK2:     [[DECL:%.+]] = getelementptr inbounds [[ST]], [[ST]]* %{{.+}}, i32 0, i32 1
     // CK2:     [[TTT:%.+]] = load double**, double*** [[DECL]],
     // CK2:     [[TTTT:%.+]] = load double*, double** [[TTT]],
     // 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]]
     // CK2:     [[CVAL:%.+]] = bitcast double*** [[CBP]] to double**
     // CK2:     [[VAL:%.+]] = load double*, double** [[CVAL]],
     // CK2:     store double* [[VAL]], double** [[PVT:%.+]],
     // CK2:     store double** [[PVT]], double*** [[PVT2:%.+]],
     // CK2:     [[TT1:%.+]] = load double**, double*** [[PVT2]],

clang/test/OpenMP/target_map_codegen.cpp

 // expected-no-diagnostics
 #ifndef HEADER
 #define HEADER
 
 ///
 /// Implicit maps.
 ///
 
 ///==========================================================================///

 // CK19: [[MTYPE08:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
 
 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK19: [[SIZE09:@.+]] = private {{.*}}constant [1 x i64] [i64 {{8|4}}]
 // CK19: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 34]
 
 // 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]
 // CK19-NOUSE: [[MTYPE16:@.+]] = private {{.*}}constant [1 x i64] [i64 33]
 
 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK19-USE: [[SIZE17:@.+]] = private {{.*}}constant [2 x i64] [i64 {{8|4}}, i64 240]
 // CK19-USE: [[MTYPE17:@.+]] = private {{.*}}constant [2 x i64] [i64 800, i64 34]

 // CK19: [[MTYPE26:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
 
 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK19: [[SIZE27:@.+]] = private {{.*}}constant [1 x i64] [i64 4]
 // CK19: [[MTYPE27:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
 
 // 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]
 // CK19-NOUSE: [[MTYPE30:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
 
 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK19-USE: [[SIZE31:@.+]] = private {{.*}}constant [4 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 {{8|4}}, i64 40]
 // CK19-USE: [[MTYPE31:@.+]] = private {{.*}}constant [4 x i64] [i64 800, i64 800, i64 800, i64 35]

 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK19-USE: [[SIZE41:@.+]] = private {{.*}}constant [3 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 208]
 // CK19-USE: [[MTYPE41:@.+]] = private {{.*}}constant [3 x i64] [i64 800, i64 800, i64 35]
 // CK19-NOUSE: [[SIZE41:@.+]] = private {{.*}}constant [1 x i64] [i64 208]
 // CK19-NOUSE: [[MTYPE41:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
 
 // 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]
 
 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK19: [[SIZE44:@.+]] = private {{.*}}constant [1 x i64] [i64 320]
 // CK19: [[MTYPE44:@.+]] = private {{.*}}constant [1 x i64] [i64 34]
 
#endif
 
   // Region 10
   // CK19-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE10]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE10]]{{.+}}, i8** null)
   // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
   // 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]]
 
   // CK19-USE: call void [[CALL10:@.+]](i32* {{[^,]+}})
   // CK19-NOUSE: call void [[CALL10:@.+]]()
   #pragma omp target map(tofrom:pa[20:60])
   {
 #ifdef USE
     pa[50]++;
 #endif
   }
 
   // Region 11
   // CK19-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE11]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE11]]{{.+}}, i8** null)
   // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
   // 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]]
 
   // CK19-USE: call void [[CALL11:@.+]](i32* {{[^,]+}})
   // CK19-NOUSE: call void [[CALL11:@.+]]()
   #pragma omp target map(alloc:pa[:60])
   {
 #ifdef USE
     pa[50]++;
 #endif
   }
 
   // Region 12
   // CK19-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE12]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE12]]{{.+}}, i8** null)
   // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
   // 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]]
 
   // CK19-USE: call void [[CALL12:@.+]](i32* {{[^,]+}})
   // CK19-NOUSE: call void [[CALL12:@.+]]()
   #pragma omp target map(to:pa[15])
   {
 #ifdef USE
     pa[15]++;
 #endif
   }
 
   // Region 13
   // CK19-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE13]]{{.+}}, i8** null)
   // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
   // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 
   // 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]]
 
   // CK19-USE: call void [[CALL13:@.+]](i32* {{[^,]+}})
   // CK19-NOUSE: call void [[CALL13:@.+]]()
   #pragma omp target map(alloc:pa[ii-23:ii])
   {
 #ifdef USE
     pa[50]++;
 #endif
   }
 
   // Region 14
   // CK19-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE14]]{{.+}}, i8** null)
   // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
   // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 
   // 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]]
 
   // CK19-USE: call void [[CALL14:@.+]](i32* {{[^,]+}})
   // CK19-NOUSE: call void [[CALL14:@.+]]()
   #pragma omp target map(to:pa[:ii])
   {
 #ifdef USE
     pa[50]++;
 #endif
   }
 
   // Region 15
   // CK19-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE15]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE15]]{{.+}}, i8** null)
   // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
   // 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]]
 
   // CK19-USE: call void [[CALL15:@.+]](i32* {{[^,]+}})
   // CK19-NOUSE: call void [[CALL15:@.+]]()
   #pragma omp target map(from:pa[ii+12])
   {
 #ifdef USE
#endif
 
   // Region 28
   // CK19-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[SIZE28]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE28]]{{.+}}, i8** null)
   // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
   // 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]],
 
   // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
   // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
   // CK19-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32****
   // CK19-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32***
   // CK19-DAG: store i32*** [[SEC0]], i32**** [[CBP1]]
#endif
 
   // Region 29
   // CK19-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[SIZE29]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE29]]{{.+}}, i8** null)
   // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
   // 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]],
 
   // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
   // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
   // CK19-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32****
   // CK19-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32***
   // CK19-DAG: store i32*** [[SEC0]], i32**** [[CBP1]]
#endif
 
   // Region 42
   // CK19-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[SIZE42]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE42]]{{.+}}, i8** null)
   // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
   // 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]],
 
   // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
   // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
   // CK19-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to double****
   // CK19-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double***
   // CK19-DAG: store double*** [[SEC0]], double**** [[CBP1]]

 // CK20: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 33]
 
 // CK20-LABEL: @.__omp_offloading_{{.*}}explicit_maps_references_and_function_args{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK20: [[SIZE02:@.+]] = private {{.*}}constant [1 x i64] [i64 4]
 // CK20: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i64] [i64 34]
 
 // 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){
 
   int &aa = a;
   float &bb = b;
   int (&cc)[10] = c;
   float *&dd = d;
void explicit_maps_references_and_function_args (int a, float b, int (&c)[10], float *d){
 
   // Region 03
   // CK20-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE03]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}, i8** null)
   // CK20-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK20-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
   // 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]]
 
   // CK20: call void [[CALL03:@.+]](float* {{[^,]+}})
   #pragma omp target map(from:d[2:3])
   {
     d[2] += 1.0f;
   }

 #ifdef CK21
 // CK21: [[ST:%.+]] = type { i32, i32, float* }
 
 // CK21-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK21: [[MTYPE00:@.+]] = private {{.*}}constant [2 x i64] [i64 32, i64 281474976710659]
 
 // 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]
 
 // CK21-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK21: [[SIZE03:@.+]] = private {{.*}}constant [1 x i64] [i64 492]
 // CK21: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i64] [i64 34]
 
#endif
 
     // Region 01
     // CK21-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}}, i8** null)
     // CK21-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
     // CK21-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
     // 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]]
 
     // CK21-USE: call void [[CALL01:@.+]](i32* {{[^,]+}})
     // CK21-NOUSE: call void [[CALL01:@.+]]()
     #pragma omp target map(lb[:X])
     {
 #ifdef USE

 // CK22: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
 
 // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK22: [[SIZE03:@.+]] = private {{.*}}constant [1 x i64] [i64 16]
 // CK22: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
 
 // 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]
 // CK22: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
 
 // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK22: [[SIZE06:@.+]] = private {{.*}}constant [1 x i64] [i64 400]
 // CK22: [[MTYPE06:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
 
 // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK22: [[SIZE07:@.+]] = private {{.*}}constant [1 x i64] [i64 {{8|4}}]
 // CK22: [[MTYPE07:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
 
 // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK22: [[SIZE08:@.+]] = private {{.*}}constant [1 x i64] [i64 16]
 // CK22: [[MTYPE08:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
 
 // 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]
 // CK22: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
 
 // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK22: [[SIZE11:@.+]] = private {{.*}}constant [1 x i64] [i64 400]
 // CK22: [[MTYPE11:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
 
 // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK22: [[SIZE12:@.+]] = private {{.*}}constant [1 x i64] [i64 {{8|4}}]
 // CK22: [[MTYPE12:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
 
 // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK22: [[SIZE13:@.+]] = private {{.*}}constant [1 x i64] [i64 16]
 // CK22: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
 
 // 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];
 int *d;
 
 struct ST {
   float fa;
 };
int explicit_maps_globals(void){
 
   // Region 04
   // CK22-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE04]]{{.+}}, i8** null)
   // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
   // 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
 
   // CK22: call void [[CALL04:@.+]](i32* {{[^,]+}})
   #pragma omp target map(d[2:5])
   { d[3]+=1; }
 
   // Region 05
int explicit_maps_globals(void){
 
   // Region 09
   // CK22-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE09]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE09]]{{.+}}, i8** null)
   // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
   // 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
 
   // CK22: call void [[CALL09:@.+]]([[ST]]* {{[^,]+}})
   #pragma omp target map(sd[2:5])
   { sd[3].fa+=1; }
 
   // Region 10
int explicit_maps_globals(void){
 
   // Region 14
   // CK22-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE14]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE14]]{{.+}}, i8** null)
   // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
   // 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
 
   // CK22: call void [[CALL14:@.+]]([[STT]]* {{[^,]+}})
   #pragma omp target map(std[2:5])
   { std[3].fa+=1; }
 
   return 0;

 // CK23: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
 
 // CK23-LABEL: @.__omp_offloading_{{.*}}explicit_maps_inside_captured{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK23: [[SIZE04:@.+]] = private {{.*}}constant [1 x i64] [i64 16]
 // CK23: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
 
 // 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){
   float b;
   float c[100];
   float *d;
 
   // CK23: call void @{{.*}}explicit_maps_inside_captured{{.*}}([[SA:%.+]]* {{.*}})
[&](void){
 
     // Region 05
     // CK23-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE05]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE05]]{{.+}}, i8** null)
     // CK23-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
     // CK23-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
     // 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:%[^,]+]]
     // CK23-DAG: [[CAP0]] = getelementptr inbounds [[SA]], [[SA]]
     // CK23-DAG: [[VAR00]] = load float**, float*** [[CAP00:%[^,]+]]
     // CK23-DAG: [[CAP00]] = getelementptr inbounds [[SA]], [[SA]]
 
     // CK23: call void [[CALL05:@.+]](float* {{[^,]+}})

 #ifdef CK27
 
 // CK27-LABEL: @.__omp_offloading_{{.*}}zero_size_section_and_private_maps{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK27: [[SIZE00:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer
 // CK27: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i64] [i64 544]
 
 // 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
 // CK27: [[SIZE05:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer
 // CK27: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i64] [i64 32]
 
 // 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
void zero_size_section_and_private_maps (int ii){
 
   // Region 01
   // CK27-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}}, i8** null)
   // CK27-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK27-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
   // 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]]
 
   // CK27: call void [[CALL01:@.+]](i32* {{[^,]+}})
   #pragma omp target map(pa[:0])
   {
     pa[50]++;
   }
 
   // Region 02
   // CK27-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}, i8** null)
   // CK27-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK27-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
   // 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]]
 
   // CK27: call void [[CALL02:@.+]](i32* {{[^,]+}})
   #pragma omp target map(pa[0:0])
   {
     pa[50]++;
   }
 
   // Region 03
   // CK27-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE03]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}, i8** null)
   // CK27-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK27-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
   // 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]]
 
   // CK27: call void [[CALL03:@.+]](i32* {{[^,]+}})
   #pragma omp target map(pa[ii:0])
   {
     pa[50]++;
   }

 #ifdef CK28
 
 // CK28-LABEL: @.__omp_offloading_{{.*}}explicit_maps_pointer_references{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
 // CK28: [[SIZE00:@.+]] = private {{.*}}constant [1 x i64] [i64 {{8|4}}]
 // CK28: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
 
 // 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){
   int *&a = p;
 
   // Region 00
   // CK28-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
   // CK28-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
void explicit_maps_pointer_references (int *p){
 
   // Region 01
   // CK28-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}}, i8** null)
   // CK28-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK28-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 
   // 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:%.+]],
   // CK28-DAG: [[VAR111]] = load i32**, i32*** [[VAR1111:%.+]],
 
   // CK28: call void [[CALL01:@.+]](i32* {{[^,]+}})
   #pragma omp target map(a[2:100])
   {

 // RUN: %clang_cc1 -DCK32 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY32 %s
 // RUN: %clang_cc1 -DCK32 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY32 %s
 // RUN: %clang_cc1 -DCK32 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY32 %s
 // SIMD-ONLY32-NOT: {{__kmpc|__tgt}}
 #ifdef 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) {
 
   // CK32-DAG: call i32 @__tgt_target_mapper(i64 -1, i8* @{{.+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE_TO]]{{.+}}, i8** null)
   // CK32-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK32-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
   // CK32-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 
   // CK32-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // 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, %{{.+}}
   // CK32-32-DAG: [[SIZE]] = sext i32 [[SZ1:%.+]] to i64
   // CK32-32-DAG: [[SZ1]] = mul nuw i32 [[SZ2:%.+]], 4
   // CK32-32-DAG: [[SZ2]] = mul nuw i32 12, %{{.+}}
   #pragma omp target map(to:([3][sa][4])f)
   f[0] = 1;
   sa = 1;
   // CK32-DAG: call i32 @__tgt_target_mapper(i64 -1, i8* @{{.+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE_FROM]]{{.+}}, i8** null)
   // CK32-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK32-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
   // CK32-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 
   // CK32-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
   // 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, %{{.+}}
   // CK32-32-DAG: [[SIZE]] = sext i32 [[SZ1:%.+]] to i64
   // CK32-32-DAG: [[SZ1]] = mul nuw i32 [[SZ2:%.+]], 5
   // CK32-32-DAG: [[SZ2]] = mul nuw i32 4, %{{.+}}
   #pragma omp target map(from: ([sa][5])f)
   f[0] = 1;
 }
 
 #endif
 #endif

clang/test/OpenMP/target_update_codegen.cpp

 // expected-no-diagnostics
 #ifndef HEADER
 #define HEADER
 
 ///==========================================================================///
 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-64
 // RUN: %clang_cc1 -DCK1 -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 CK1 --check-prefix CK1-64
 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK1 --check-prefix CK1-32

 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // RUN: %clang_cc1 -DCK5 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // RUN: %clang_cc1 -DCK5 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
 
 #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]]
   // CK5-DAG: store i32* [[B_VAL_2:%.+]], i32** [[PC0]]
   // CK5-DAG: [[B_VAL_2]] = load i32*, i32** [[B_ADDR]]
   #pragma omp target update to(*B)
   *B += e;
   #pragma omp target update from(*B)

 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // RUN: %clang_cc1 -DCK6 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // RUN: %clang_cc1 -DCK6 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
 
 #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]]
   // CK6-DAG: store i32* [[ADD_PTR:%.+]], i32** [[PC0]]
   // CK6-64-DAG: [[ADD_PTR]] = getelementptr inbounds i32, i32* [[ONE:%.+]], i{{32|64}} [[IDX_EXT:%.+]]
   // CK6-32-DAG: [[ADD_PTR]] = getelementptr inbounds i32, i32* [[ONE:%.+]], i{{32|64}} [[L_VAL:%.+]]
   // CK6-64-DAG: [[IDX_EXT]] = sext i32 [[L_VAL:%.+]] to i64
   // CK6-DAG: [[L_VAL]] = load i32, i32* [[L_ADDR:%.+]]

 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // RUN: %clang_cc1 -DCK7 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // RUN: %clang_cc1 -DCK7 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
 
 #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]]
   // CK7-DAG: store i32* [[ARRAY_IDX:%.+]], i32** [[PC0]]
   // CK7-DAG: [[ARRAY_IDX]] = getelementptr inbounds i32, i32* [[ADD_PTR:%.+]], i{{32|64}} [[IDX_PROM:%.+]]
   // CK7-64-DAG: [[ADD_PTR]] = getelementptr inbounds i32, i32* [[ONE:%.+]], i64 [[IDX_EXT:%.+]]
   // CK7-32-DAG: [[ADD_PTR]] = getelementptr inbounds i32, i32* [[B_VAL_2:%.+]], i32 [[L_VAL:%.+]]
   // CK7-32-DAG: [[B_VAL_2]] = load i32*, i32** [[B_ADDR]]

 // RUN: %clang_cc1 -DCK8 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // RUN: %clang_cc1 -DCK8 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // RUN: %clang_cc1 -DCK8 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // 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]]
   // CK8-DAG: store i32* [[ARRAY_IDX_4:%.+]], i32** [[PC0]]
   // CK8-DAG: store i32** [[ARRAY_IDX_1]], i32*** [[NINE:%.+]]
   // 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])
 }
 
 #endif
 ///==========================================================================///
 // RUN: %clang_cc1 -DCK9 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK9 --check-prefix CK9-64

 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
 
 #ifdef CK9
 
 struct S {
   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]]
   // CK9-DAG: store double* [[P_VAL:%.+]], double** [[PC0]]
   // CK9-DAG: store i{{.+}} 8, i{{.+}}* [[SIZE0]]
   // CK9-DAG: [[P]] = getelementptr inbounds [[STRUCT_S:%.+]], [[STRUCT_S]]* [[S_VAL:%.+]], i32 0, i32 0
   // CK9-DAG: [[S_VAL]] = load [[STRUCT_S]]*, [[STRUCT_S]]** [[S_ADDR:%.+]]
   // CK9-DAG: [[P_VAL]] = load double*, double** [[P_1:%.+]],

 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
 #ifdef CK10
 
 struct S {
   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]]
   // CK10-DAG: store double* [[ADD_PTR:%.+]], double** [[PC0]]
   // CK10-DAG: store i{{.+}} 8, i{{.+}}* [[SIZE0]]
   // CK10-64-DAG: [[ADD_PTR]] = getelementptr inbounds double, double* [[S_P:%.+]], i{{.+}} [[IDX_EXT:%.+]]
   // CK10-32-DAG: [[ADD_PTR]] = getelementptr inbounds double, double* [[S_P:%.+]], i{{.+}} [[L_VAL:%.+]]
   // CK10-64-DAG: [[IDX_EXT]] = sext i32 [[L_VAL:%.+]] to i64

 // RUN: %clang_cc1 -DCK11 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
 #ifdef CK11
 
 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]]
   // CK11-DAG: store double* [[ARRAY_IDX:%.+]], double** [[PC0]]
   // CK11-DAG: store i{{.+}} 8, i{{.+}} [[SIZE0]]
   // CK11-DAG: [[P]] = getelementptr inbounds [[STRUCT_S:%.+]], [[STRUCT_S]]* [[SS_1:%.+]], i32 0, i32 0
   // CK11-DAG: [[ARRAY_IDX]] = getelementptr inbounds double, double* [[ADD_PTR:%.+]], i{{.+}} 3
   // CK11-64-DAG: [[ADD_PTR]] = getelementptr inbounds double, double* [[S_P:%.+]], i{{.+}} [[IDX_EXT:%.+]]

 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
 #ifdef CK12
 
 struct S {
   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))
 }
 
 #endif
 ///==========================================================================///
 // RUN: %clang_cc1 -DCK13 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK13 --check-prefix CK13-64
 // RUN: %clang_cc1 -DCK13 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK13 --check-prefix CK13-64
 // RUN: %clang_cc1 -DCK13 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK13 --check-prefix CK13-32
 // RUN: %clang_cc1 -DCK13 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK13 --check-prefix CK13-32
 
 // RUN: %clang_cc1 -DCK13 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // RUN: %clang_cc1 -DCK13 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // RUN: %clang_cc1 -DCK13 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // RUN: %clang_cc1 -DCK13 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // 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]]
   // CK13-DAG: store i32* [[ADD_PTR_2:%.+]], i32** [[PC0]]
   // CK13-64-DAG: [[ADD_PTR_2]] = getelementptr inbounds i32, i32* [[RESULT:%.+]], i64 [[IDX_EXT_1:%.+]]
   // CK13-32-DAG: [[ADD_PTR_2]] = getelementptr inbounds i32, i32* [[RESULT:%.+]], i32 [[B_ADDR:%.+]]
   // CK13-64-DAG: [[IDX_EXT_1]] = sext i32 [[B_ADDR:%.+]]
   // CK13-DAG: [[RESULT]] = load i32*, i32** [[ADD_PTR:%.+]],

 // RUN: %clang_cc1 -DCK15 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // RUN: %clang_cc1 -DCK15 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // RUN: %clang_cc1 -DCK15 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // 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;
   double *≺
   SSA(double *&pr) : pr(pr) {}
 };
 
 //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]]
   // CK15-DAG: store double* [[ADD_PTR:%.+]], double** [[PC2]]
   // CK15-DAG: store i64 8, i64* [[SIZE2]]
   // 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]]
   // CK15-DAG: [[COMPUTE_SIZE]] = sdiv exact i64 [[NINE:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
   // CK15-DAG: [[NINE]] = sub i{{.+}} [[SEVEN:%.+]], [[EIGHT:%.+]]
   // CK15-DAG: [[SEVEN]] = ptrtoint i8* [[SIX:%.+]] to i64
   // CK15-DAG: [[EIGHT]] = ptrtoint i8* [[FIVE:%.+]] to i64
   // CK15-DAG: [[SIX]] = bitcast double** {{.+}} to i8*

 // RUN: %clang_cc1 -DCK16 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // RUN: %clang_cc1 -DCK16 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // RUN: %clang_cc1 -DCK16 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // 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]]
   // CK16-DAG: store float* [[ADD_PTR:%.+]], float** [[PC0]]
   // CK16-32-DAG: [[ADD_PTR]] = getelementptr inbounds float, float* [[THREE:%.+]], i32 [[I:%.+]]
   // CK16-64-DAG: [[ADD_PTR]] = getelementptr inbounds float, float* [[THREE:%.+]], i64 [[IDX_EXT:%.+]]
   // CK16-DAG: [[THREE]] = load float*, float** [[F_ADDR]],
   // CK16-64-DAG: [[IDX_EXT]] = sext i32 [[I:%.+]] to i64

 // RUN: %clang_cc1 -DCK17 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // RUN: %clang_cc1 -DCK17 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // RUN: %clang_cc1 -DCK17 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 // 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;
   SSA *sa;
 };
 
 //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]],
   // CK17-DAG: store float* [[ADD_PTR_4:%.+]], float** [[PC0]],
   // CK17-64-DAG: [[ADD_PTR_4]] = getelementptr inbounds float, float* [[SEVEN:%.+]], i64 [[IDX_EXT_3:%.+]]
   // CK17-64-DAG: [[IDX_EXT_3]] = sext i32 [[I_VAL:%.+]] to i64
   // CK17-32-DAG: [[ADD_PTR_4]] = getelementptr inbounds float, float* [[SEVEN:%.+]], i32 [[I_VAL:%.+]]
   // CK17-DAG: [[I_VAL]] = load i32, i32* [[I:%.+]],

 // RUN: %clang_cc1 -DCK18 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY18 %s
 // RUN: %clang_cc1 -DCK18 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY18 %s
 // RUN: %clang_cc1 -DCK18 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY18 %s
 // 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]],
 
   // CK18-64-DAG: [[SIZE]] = mul nuw i64 [[SZ1:%.+]], 4
   // CK18-64-DAG: [[SZ1]] = mul nuw i64 12, %{{.+}}
   // CK18-32-DAG: [[SIZE]] = sext i32 [[SZ1:%.+]] to i64
   // CK18-32-DAG: [[SZ1]] = mul nuw i32 [[SZ2:%.+]], 4
   // 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]],
 
   // CK18-64-DAG: [[SIZE]] = mul nuw i64 [[SZ1:%.+]], 5
   // CK18-64-DAG: [[SZ1]] = mul nuw i64 4, %{{.+}}
   // CK18-32-DAG: [[SIZE]] = sext i32 [[SZ1:%.+]] to i64
   // CK18-32-DAG: [[SZ1]] = mul nuw i32 [[SZ2:%.+]], 5
   // CK18-32-DAG: [[SZ2]] = mul nuw i32 4, %{{.+}}
   #pragma omp target update from(([sa][5])f)
 }
 
 #endif
 #endif

openmp/libomptarget/src/omptarget.cpp

#endif
       if (arg_types[i] & OMP_TGT_MAPTYPE_TO) {
         int rt = Device.data_submit(TgtPtrBegin, HstPtrBegin, arg_sizes[i],
                                     &AsyncInfo);
         if (rt != OFFLOAD_SUCCESS) {
           DP("Copying data to device failed, failed.\n");
           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;
 #ifdef OMPTARGET_DEBUG
       void *TgtPtrBase = (void *)((intptr_t)TgtPtrBegin + TgtBaseOffset);
       DP("Obtained target argument " DPxMOD " from host pointer " DPxMOD "\n",
           DPxPTR(TgtPtrBase), DPxPTR(HstPtrBegin));
 #endif