From: Samuel Antao <sfantao@us.ibm.com>
Date: Thu, 26 May 2016 16:48:10 +0000 (+0000)
Subject: [OpenMP] Adjust map type bits according to latest spec and use zero size array sectio... 
X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=42720cb0acda06dd20caf809a0c4d931ad249dcf;p=clang

[OpenMP] Adjust map type bits according to latest spec and use zero size array sections for pointers.

Summary: This patch changes the bits used to specify the map types according to the latest version of the libomptarget document and add the support for zero size array section when pointers are being implicitly mapped. This completes the missing new 4.5 map semantics.

Reviewers: hfinkel, carlo.bertolli, arpith-jacob, kkwli0, ABataev

Subscribers: caomhin, cfe-commits

Differential Revision: http://reviews.llvm.org/D20111

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@270868 91177308-0d34-0410-b5e6-96231b3b80d8
---

diff --git a/lib/CodeGen/CGOpenMPRuntime.cpp b/lib/CodeGen/CGOpenMPRuntime.cpp
index 4b0d21373e..3ac572e8be 100644
--- a/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -4909,8 +4909,6 @@ public:
   /// \brief Values for bit flags used to specify the mapping type for
   /// offloading.
   enum OpenMPOffloadMappingFlags {
-    /// \brief Only allocate memory on the device,
-    OMP_MAP_ALLOC = 0x00,
     /// \brief Allocate memory on the device and move data from host to device.
     OMP_MAP_TO = 0x01,
     /// \brief Allocate memory on the device and move data from device to host.
@@ -4918,19 +4916,19 @@ public:
     /// \brief Always perform the requested mapping action on the element, even
     /// if it was already mapped before.
     OMP_MAP_ALWAYS = 0x04,
-    /// \brief Decrement the reference count associated with the element without
-    /// executing any other action.
-    OMP_MAP_RELEASE = 0x08,
     /// \brief Delete the element from the device environment, ignoring the
     /// current reference count associated with the element.
-    OMP_MAP_DELETE = 0x10,
-    /// \brief The element passed to the device is a pointer.
-    OMP_MAP_PTR = 0x20,
-    /// \brief Signal the element as extra, i.e. is not argument to the target
-    /// region kernel.
-    OMP_MAP_EXTRA = 0x40,
+    OMP_MAP_DELETE = 0x08,
+    /// \brief The element being mapped is a pointer, therefore the pointee
+    /// should be mapped as well.
+    OMP_MAP_IS_PTR = 0x10,
+    /// \brief This flags signals that an argument is the first one relating to
+    /// a map/private clause expression. For some cases a single
+    /// map/privatization results in multiple arguments passed to the runtime
+    /// library.
+    OMP_MAP_FIRST_REF = 0x20,
     /// \brief Pass the element to the device by value.
-    OMP_MAP_BYCOPY = 0x80,
+    OMP_MAP_PRIVATE_VAL = 0x100,
   };
 
   typedef SmallVector<llvm::Value *, 16> MapValuesArrayTy;
@@ -4987,14 +4985,19 @@ private:
 
   /// \brief Return the corresponding bits for a given map clause modifier. Add
   /// a flag marking the map as a pointer if requested. Add a flag marking the
-  /// map as extra, meaning is not an argument of the kernel.
+  /// map as the first one of a series of maps that relate to the same map
+  /// expression.
   unsigned getMapTypeBits(OpenMPMapClauseKind MapType,
                           OpenMPMapClauseKind MapTypeModifier, bool AddPtrFlag,
-                          bool AddExtraFlag) const {
+                          bool AddIsFirstFlag) const {
     unsigned Bits = 0u;
     switch (MapType) {
     case OMPC_MAP_alloc:
-      Bits = OMP_MAP_ALLOC;
+    case OMPC_MAP_release:
+      // alloc and release is the default behavior in the runtime library,  i.e.
+      // if we don't pass any bits alloc/release that is what the runtime is
+      // going to do. Therefore, we don't need to signal anything for these two
+      // type modifiers.
       break;
     case OMPC_MAP_to:
       Bits = OMP_MAP_TO;
@@ -5008,17 +5011,14 @@ private:
     case OMPC_MAP_delete:
       Bits = OMP_MAP_DELETE;
       break;
-    case OMPC_MAP_release:
-      Bits = OMP_MAP_RELEASE;
-      break;
     default:
       llvm_unreachable("Unexpected map type!");
       break;
     }
     if (AddPtrFlag)
-      Bits |= OMP_MAP_PTR;
-    if (AddExtraFlag)
-      Bits |= OMP_MAP_EXTRA;
+      Bits |= OMP_MAP_IS_PTR;
+    if (AddIsFirstFlag)
+      Bits |= OMP_MAP_FIRST_REF;
     if (MapTypeModifier == OMPC_MAP_always)
       Bits |= OMP_MAP_ALWAYS;
     return Bits;
@@ -5270,13 +5270,13 @@ private:
 
         Pointers.push_back(LB);
         Sizes.push_back(Size);
-        // We need to add a pointer flag for each map that comes from the the
-        // same expression except for the first one. We need to add the extra
-        // flag for each map that relates with the current capture, except for
-        // the first one (there is a set of entries for each capture).
+        // 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).
         Types.push_back(getMapTypeBits(MapType, MapTypeModifier,
                                        !IsExpressionFirstInfo,
-                                       !IsCaptureFirstInfo));
+                                       IsCaptureFirstInfo));
 
         // If we have a final array section, we are done with this expression.
         if (IsFinalArraySection)
@@ -5583,7 +5583,8 @@ void CGOpenMPRuntime::emitTargetCall(CodeGenFunction &CGF,
       CurPointers.push_back(*CV);
       CurSizes.push_back(CGF.getTypeSize(RI->getType()));
       // Copy to the device as an argument. No need to retrieve it.
-      CurMapTypes.push_back(MappableExprsHandler::OMP_MAP_BYCOPY);
+      CurMapTypes.push_back(MappableExprsHandler::OMP_MAP_PRIVATE_VAL |
+                            MappableExprsHandler::OMP_MAP_FIRST_REF);
     } else {
       // If we have any information in the map clause, we use it, otherwise we
       // just do a default mapping.
@@ -5602,10 +5603,10 @@ void CGOpenMPRuntime::emitTargetCall(CodeGenFunction &CGF,
           CurMapTypes.push_back(MappableExprsHandler::OMP_MAP_TO |
                                 MappableExprsHandler::OMP_MAP_FROM);
         } else if (CI->capturesVariableByCopy()) {
-          CurMapTypes.push_back(MappableExprsHandler::OMP_MAP_BYCOPY);
           if (!RI->getType()->isAnyPointerType()) {
             // If the field is not a pointer, we need to save the actual value
             // and load it as a void pointer.
+            CurMapTypes.push_back(MappableExprsHandler::OMP_MAP_PRIVATE_VAL);
             auto DstAddr = CGF.CreateMemTemp(
                 Ctx.getUIntPtrType(),
                 Twine(CI->getCapturedVar()->getName()) + ".casted");
@@ -5624,11 +5625,17 @@ void CGOpenMPRuntime::emitTargetCall(CodeGenFunction &CGF,
             CurBasePointers.push_back(
                 CGF.EmitLoadOfLValue(DstLV, SourceLocation()).getScalarVal());
             CurPointers.push_back(CurBasePointers.back());
+
+            // Get the size of the type to be used in the map.
+            CurSizes.push_back(CGF.getTypeSize(RI->getType()));
           } else {
+            // Pointers are implicitly mapped with a zero size and no flags
+            // (other than first map that is added for all implicit maps).
+            CurMapTypes.push_back(0u);
             CurBasePointers.push_back(*CV);
             CurPointers.push_back(*CV);
+            CurSizes.push_back(llvm::Constant::getNullValue(CGM.SizeTy));
           }
-          CurSizes.push_back(CGF.getTypeSize(RI->getType()));
         } else {
           assert(CI->capturesVariable() && "Expected captured reference.");
           CurBasePointers.push_back(*CV);
@@ -5640,13 +5647,15 @@ void CGOpenMPRuntime::emitTargetCall(CodeGenFunction &CGF,
           CurSizes.push_back(CGF.getTypeSize(ElementType));
           // The default map type for a scalar/complex type is 'to' because by
           // default the value doesn't have to be retrieved. For an aggregate
-          // type,
-          // the default is 'tofrom'.
+          // type, the default is 'tofrom'.
           CurMapTypes.push_back(ElementType->isAggregateType()
                                     ? (MappableExprsHandler::OMP_MAP_TO |
                                        MappableExprsHandler::OMP_MAP_FROM)
                                     : MappableExprsHandler::OMP_MAP_TO);
         }
+        // Every default map produces a single argument, so, it is always the
+        // first one.
+        CurMapTypes.back() |= MappableExprsHandler::OMP_MAP_FIRST_REF;
       }
     }
     // We expect to have at least an element of information for this capture.
diff --git a/test/OpenMP/target_codegen.cpp b/test/OpenMP/target_codegen.cpp
index 6f38bc6d7a..525d26e158 100644
--- a/test/OpenMP/target_codegen.cpp
+++ b/test/OpenMP/target_codegen.cpp
@@ -33,15 +33,15 @@
 // sizes.
 
 // CHECK-DAG: [[SIZET2:@.+]] = private unnamed_addr constant [1 x i{{32|64}}] [i[[SZ:32|64]] 2]
-// CHECK-DAG: [[MAPT2:@.+]] = private unnamed_addr constant [1 x i32] [i32 128]
+// CHECK-DAG: [[MAPT2:@.+]] = private unnamed_addr constant [1 x i32] [i32 288]
 // CHECK-DAG: [[SIZET3:@.+]] = private unnamed_addr constant [2 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2]
-// CHECK-DAG: [[MAPT3:@.+]] = private unnamed_addr constant [2 x i32] [i32 128, i32 128]
-// CHECK-DAG: [[MAPT4:@.+]] = private unnamed_addr constant [9 x i32] [i32 128, i32 3, i32 128, i32 3, i32 3, i32 128, i32 128, i32 3, i32 3]
+// CHECK-DAG: [[MAPT3:@.+]] = private unnamed_addr constant [2 x i32] [i32 288, i32 288]
+// CHECK-DAG: [[MAPT4:@.+]] = private unnamed_addr constant [9 x i32] [i32 288, i32 35, i32 288, i32 35, i32 35, i32 288, i32 288, i32 35, i32 35]
 // CHECK-DAG: [[SIZET5:@.+]] = private unnamed_addr constant [3 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2, i[[SZ]] 40]
-// CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [3 x i32] [i32 128, i32 128, i32 3]
+// CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [3 x i32] [i32 288, i32 288, i32 35]
 // CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [4 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2, i[[SZ]] 1, i[[SZ]] 40]
-// CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [4 x i32] [i32 128, i32 128, i32 128, i32 3]
-// CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [5 x i32] [i32 3, i32 128, i32 128, i32 128, i32 3]
+// CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [4 x i32] [i32 288, i32 288, i32 288, i32 35]
+// CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [5 x i32] [i32 35, i32 288, i32 288, i32 288, i32 35]
 // CHECK-DAG: @{{.*}} = private constant i8 0
 // CHECK-DAG: @{{.*}} = private constant i8 0
 // CHECK-DAG: @{{.*}} = private constant i8 0
diff --git a/test/OpenMP/target_codegen_registration.cpp b/test/OpenMP/target_codegen_registration.cpp
index 3afdad9153..c26752dd78 100644
--- a/test/OpenMP/target_codegen_registration.cpp
+++ b/test/OpenMP/target_codegen_registration.cpp
@@ -61,40 +61,40 @@
 // CHECK-DAG: {{@.+}} = private constant i8 0
 // TCHECK-NOT: {{@.+}} = private constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
-// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 128]
+// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
 // CHECK-DAG: {{@.+}} = private constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
-// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 128]
+// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
 // CHECK-DAG: {{@.+}} = private constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
-// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 128]
+// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
 // CHECK-DAG: {{@.+}} = private constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
-// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 128]
+// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
 // CHECK-DAG: {{@.+}} = private constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
-// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 128]
+// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
 // CHECK-DAG: {{@.+}} = private constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
-// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 128]
+// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
 // CHECK-DAG: {{@.+}} = private constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
-// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 128]
+// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
 // CHECK-DAG: {{@.+}} = private constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
-// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 128]
+// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
 // CHECK-DAG: {{@.+}} = private constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
-// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 128]
+// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
 // CHECK-DAG: {{@.+}} = private constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
-// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 128]
+// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
 // CHECK-DAG: {{@.+}} = private constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
-// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 128]
+// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
 // CHECK-DAG: {{@.+}} = private constant i8 0
 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
-// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 128]
+// CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
 
 // CHECK-NTARGET-NOT: private constant i8 0
 // CHECK-NTARGET-NOT: private unnamed_addr constant [1 x i
diff --git a/test/OpenMP/target_data_codegen.cpp b/test/OpenMP/target_data_codegen.cpp
index c8b3aca37c..e2433ffc76 100644
--- a/test/OpenMP/target_data_codegen.cpp
+++ b/test/OpenMP/target_data_codegen.cpp
@@ -22,15 +22,15 @@ ST<int> gb;
 double gc[100];
 
 // CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 800]
-// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i32] [i32 2]
+// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i32] [i32 34]
 
 // CK1: [[SIZE02:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 4]
-// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i32] [i32 1]
+// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i32] [i32 33]
 
-// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i32] [i32 5]
+// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i32] [i32 37]
 
 // CK1: [[SIZE04:@.+]] = {{.+}}constant [2 x i[[sz]]] [i[[sz]] {{8|4}}, i[[sz]] 24]
-// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i32] [i32 1, i32 97]
+// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i32] [i32 33, i32 17]
 
 // CK1-LABEL: _Z3fooi
 void foo(int arg) {
@@ -173,7 +173,7 @@ struct ST {
 };
 
 // CK2: [[SIZE00:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}, i{{64|32}} 24]
-// CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i32] [i32 5, i32 101]
+// CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i32] [i32 37, i32 21]
 
 // CK2-LABEL: _Z3bari
 int bar(int arg){
diff --git a/test/OpenMP/target_enter_data_codegen.cpp b/test/OpenMP/target_enter_data_codegen.cpp
index 595f70a0c2..c40a7c50bb 100644
--- a/test/OpenMP/target_enter_data_codegen.cpp
+++ b/test/OpenMP/target_enter_data_codegen.cpp
@@ -22,15 +22,15 @@ ST<int> gb;
 double gc[100];
 
 // CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 800]
-// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i32] zeroinitializer
+// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i32] [i32 32]
 
 // CK1: [[SIZE02:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 4]
-// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i32] [i32 1]
+// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i32] [i32 33]
 
-// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i32] [i32 5]
+// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i32] [i32 37]
 
 // CK1: [[SIZE04:@.+]] = {{.+}}constant [2 x i[[sz]]] [i[[sz]] {{8|4}}, i[[sz]] 24]
-// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i32] [i32 1, i32 97]
+// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i32] [i32 33, i32 17]
 
 // CK1-LABEL: _Z3fooi
 void foo(int arg) {
@@ -156,7 +156,7 @@ struct ST {
 };
 
 // CK2: [[SIZE00:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}, i{{64|32}} 24]
-// CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i32] [i32 5, i32 101]
+// CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i32] [i32 37, i32 21]
 
 // CK2-LABEL: _Z3bari
 int bar(int arg){
diff --git a/test/OpenMP/target_exit_data_codegen.cpp b/test/OpenMP/target_exit_data_codegen.cpp
index 52d096c4ea..15d3ec1b09 100644
--- a/test/OpenMP/target_exit_data_codegen.cpp
+++ b/test/OpenMP/target_exit_data_codegen.cpp
@@ -22,15 +22,15 @@ ST<int> gb;
 double gc[100];
 
 // CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 800]
-// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i32] [i32 2]
+// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i32] [i32 34]
 
 // CK1: [[SIZE02:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 4]
-// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i32] [i32 8]
+// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i32] [i32 32]
 
-// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i32] [i32 6]
+// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i32] [i32 38]
 
 // CK1: [[SIZE04:@.+]] = {{.+}}constant [2 x i[[sz]]] [i[[sz]] {{8|4}}, i[[sz]] 24]
-// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i32] [i32 8, i32 104]
+// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i32] [i32 32, i32 16]
 
 // CK1-LABEL: _Z3fooi
 void foo(int arg) {
@@ -156,7 +156,7 @@ struct ST {
 };
 
 // CK2: [[SIZE00:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}, i{{64|32}} 24]
-// CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i32] [i32 12, i32 108]
+// CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i32] [i32 36, i32 20]
 
 // CK2-LABEL: _Z3bari
 int bar(int arg){
diff --git a/test/OpenMP/target_firstprivate_codegen.cpp b/test/OpenMP/target_firstprivate_codegen.cpp
index a3e2b9a3fe..05f1625f2b 100644
--- a/test/OpenMP/target_firstprivate_codegen.cpp
+++ b/test/OpenMP/target_firstprivate_codegen.cpp
@@ -33,16 +33,15 @@ struct TT{
 // TCHECK:  [[S1:%.+]] = type { double }
 
 // CHECK-DAG:  [[SIZET:@.+]] = private unnamed_addr constant [1 x i{{32|64}}] [i[[SZ:32|64]] 4]
-// CHECK:  [[MAPT:@.+]] = private unnamed_addr constant [1 x i32] [i32 128]
-// CHECK-DAG:  [[MAPT2:@.+]] = private unnamed_addr constant [9 x i32] [i32 128, i32 3, i32 128, i32 3, i32 3, i32 128, i32 128, i32 3, i32 3]
-// CHECK-64-DAG:  [[SIZET3:@.+]] = private unnamed_addr constant [1 x i{{32|64}}] [i[[SZ]] 8]
-// CHECK-32-DAG:  [[SIZET3:@.+]] = private unnamed_addr constant [1 x i32] [i[[SZ]] 4]
-// CHECK-DAG:  [[MAPT3:@.+]] = private unnamed_addr constant [1 x i32] [i32 128]
-// CHECK-DAG:  [[MAPT4:@.+]] = private unnamed_addr constant [5 x i32] [i32 3, i32 128, i32 128, i32 128, i32 3]
+// CHECK:  [[MAPT:@.+]] = private unnamed_addr constant [1 x i32] [i32 288]
+// CHECK-DAG:  [[MAPT2:@.+]] = private unnamed_addr constant [9 x i32] [i32 288, i32 35, i32 288, i32 35, i32 35, i32 288, i32 288, i32 35, i32 35]
+// CHECK-DAG:  [[SIZET3:@.+]] = private unnamed_addr constant [1 x i{{32|64}}] zeroinitializer
+// CHECK-DAG:  [[MAPT3:@.+]] = private unnamed_addr constant [1 x i32] [i32 32]
+// CHECK-DAG:  [[MAPT4:@.+]] = private unnamed_addr constant [5 x i32] [i32 35, i32 288, i32 288, i32 288, i32 35]
 // CHECK-DAG:  [[SIZET5:@.+]] = private unnamed_addr constant [3 x i{{32|64}}] [i[[SZ]] 4, i[[SZ]] 1, i[[SZ]] 40]
-// CHECK-DAG:  [[MAPT5:@.+]] = private unnamed_addr constant [3 x i32] [i32 128, i32 128, i32 3]
+// CHECK-DAG:  [[MAPT5:@.+]] = private unnamed_addr constant [3 x i32] [i32 288, i32 288, i32 35]
 // CHECK-DAG:  [[SIZET6:@.+]] = private unnamed_addr constant [2 x i{{32|64}}] [i[[SZ]] 4, i[[SZ]] 40]
-// CHECK-DAG:  [[MAPT6:@.+]] = private unnamed_addr constant [2 x i32] [i32 128, i32 3]
+// CHECK-DAG:  [[MAPT6:@.+]] = private unnamed_addr constant [2 x i32] [i32 288, i32 35]
 
 
 // CHECK: define {{.*}}[[FOO:@.+]](
diff --git a/test/OpenMP/target_map_codegen.cpp b/test/OpenMP/target_map_codegen.cpp
index 03a514d11f..2f6cf3ccea 100644
--- a/test/OpenMP/target_map_codegen.cpp
+++ b/test/OpenMP/target_map_codegen.cpp
@@ -16,8 +16,8 @@
 #ifdef CK1
 
 // CK1-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
-// Map types: OMP_MAP_BYCOPY = 128
-// CK1-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
+// Map types: OMP_MAP_PRIVATE_VAL | OMP_MAP_IS_FIRST = 288
+// CK1-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 288]
 
 // CK1-LABEL: implicit_maps_integer
 void implicit_maps_integer (int a){
@@ -61,8 +61,8 @@ void implicit_maps_integer (int a){
 #ifdef CK2
 
 // CK2-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
-// Map types: OMP_MAP_BYCOPY = 128
-// CK2-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
+// Map types: OMP_MAP_PRIVATE_VAL | OMP_MAP_IS_FIRST = 288
+// CK2-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 288]
 
 // CK2-LABEL: implicit_maps_integer_reference
 void implicit_maps_integer_reference (int a){
@@ -110,8 +110,8 @@ void implicit_maps_integer_reference (int a){
 #ifdef CK3
 
 // CK3-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
-// Map types: OMP_MAP_BYCOPY = 128
-// CK3-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
+// Map types: OMP_MAP_PRIVATE_VAL | OMP_MAP_IS_FIRST = 288
+// CK3-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 288]
 
 // CK3-LABEL: implicit_maps_parameter
 void implicit_maps_parameter (int a){
@@ -154,8 +154,8 @@ void implicit_maps_parameter (int a){
 #ifdef CK4
 
 // CK4-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
-// Map types: OMP_MAP_BYCOPY = 128
-// CK4-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
+// Map types: OMP_MAP_PRIVATE_VAL | OMP_MAP_IS_FIRST = 288
+// CK4-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 288]
 
 // CK4-LABEL: implicit_maps_nested_integer
 void implicit_maps_nested_integer (int a){
@@ -210,8 +210,8 @@ void implicit_maps_nested_integer (int a){
 #ifdef CK5
 
 // CK5-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
-// Map types: OMP_MAP_BYCOPY = 128
-// CK5-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
+// Map types: OMP_MAP_PRIVATE_VAL | OMP_MAP_IS_FIRST = 288
+// CK5-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 288]
 
 // CK5-LABEL: implicit_maps_nested_integer_and_enum
 void implicit_maps_nested_integer_and_enum (int a){
@@ -261,8 +261,8 @@ void implicit_maps_nested_integer_and_enum (int a){
 #ifdef CK6
 // CK6-DAG: [[GBL:@Gi]] = global i32 0
 // CK6-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
-// Map types: OMP_MAP_BYCOPY = 128
-// CK6-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
+// Map types: OMP_MAP_PRIVATE_VAL | OMP_MAP_IS_FIRST = 288
+// CK6-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 288]
 
 // CK6-LABEL: implicit_maps_host_global
 int Gi;
@@ -310,10 +310,10 @@ void implicit_maps_host_global (int a){
 // therefore it is passed by reference with a map 'to' specification.
 
 // CK7-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 8]
-// Map types: OMP_MAP_BYCOPY = 128
-// CK7-64-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
-// Map types: OMP_MAP_TO = 1
-// CK7-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 1]
+// Map types: OMP_MAP_PRIVATE_VAL | OMP_MAP_IS_FIRST = 288
+// CK7-64-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 288]
+// Map types: OMP_MAP_TO  | OMP_MAP_IS_FIRST = 33
+// CK7-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 33]
 
 // CK7-LABEL: implicit_maps_double
 void implicit_maps_double (int a){
@@ -369,8 +369,8 @@ void implicit_maps_double (int a){
 #ifdef CK8
 
 // CK8-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
-// Map types: OMP_MAP_BYCOPY = 128
-// CK8-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
+// Map types: OMP_MAP_PRIVATE_VAL | OMP_MAP_IS_FIRST = 288
+// CK8-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 288]
 
 // CK8-LABEL: implicit_maps_float
 void implicit_maps_float (int a){
@@ -413,8 +413,8 @@ void implicit_maps_float (int a){
 #ifdef CK9
 
 // CK9-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 16]
-// Map types: OMP_MAP_TO + OMP_MAP_FROM = 2 + 1
-// CK9-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 3]
+// Map types: OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST = 35
+// CK9-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 35]
 
 // CK9-LABEL: implicit_maps_array
 void implicit_maps_array (int a){
@@ -453,9 +453,9 @@ void implicit_maps_array (int a){
 // RUN: %clang_cc1 -fopenmp -fomptargets=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 CK10
 #ifdef CK10
 
-// CK10-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}]
-// Map types: OMP_MAP_BYCOPY = 128
-// CK10-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
+// CK10-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] zeroinitializer
+// Map types: OMP_MAP_IS_FIRST = 32
+// CK10-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 32]
 
 // CK10-LABEL: implicit_maps_pointer
 void implicit_maps_pointer (){
@@ -496,8 +496,8 @@ void implicit_maps_pointer (){
 #ifdef CK11
 
 // CK11-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 16]
-// Map types: OMP_MAP_TO = 1
-// CK11-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 1]
+// Map types: OMP_MAP_TO + OMP_MAP_IS_FIRST = 33
+// CK11-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 33]
 
 // CK11-LABEL: implicit_maps_double_complex
 void implicit_maps_double_complex (int a){
@@ -539,10 +539,10 @@ void implicit_maps_double_complex (int a){
 // therefore it is passed by reference with a map 'to' specification.
 
 // CK12-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 8]
-// Map types: OMP_MAP_BYCOPY = 128
-// CK12-64-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
-// Map types: OMP_MAP_TO = 1
-// CK12-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 1]
+// Map types: OMP_MAP_PRIVATE_VAL + OMP_MAP_IS_FIRST = 288
+// CK12-64-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 288]
+// Map types: OMP_MAP_TO + OMP_MAP_IS_FIRST = 33
+// CK12-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 33]
 
 // CK12-LABEL: implicit_maps_float_complex
 void implicit_maps_float_complex (int a){
@@ -598,10 +598,10 @@ void implicit_maps_float_complex (int a){
 
 // We don't have a constant map size for VLAs.
 // Map types:
-//  - OMP_MAP_BYCOPY = 128 (vla size)
-//  - OMP_MAP_BYCOPY = 128 (vla size)
-//  - OMP_MAP_TO + OMP_MAP_FROM = 2 + 1
-// CK13-DAG: [[TYPES:@.+]] = {{.+}}constant [3 x i32] [i32 128, i32 128, i32 3]
+//  - OMP_MAP_PRIVATE_VAL + OMP_MAP_IS_FIRST = 288 (vla size)
+//  - OMP_MAP_PRIVATE_VAL + OMP_MAP_IS_FIRST = 288 (vla size)
+//  - OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST = 35
+// CK13-DAG: [[TYPES:@.+]] = {{.+}}constant [3 x i32] [i32 288, i32 288, i32 35]
 
 // CK13-LABEL: implicit_maps_variable_length_array
 void implicit_maps_variable_length_array (int a){
@@ -669,9 +669,9 @@ void implicit_maps_variable_length_array (int a){
 // CK14-DAG: [[ST:%.+]] = type { i32, double }
 // CK14-DAG: [[SIZES:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{16|12}}, i{{64|32}} 4]
 // Map types:
-// - OMP_MAP_TO | OMP_MAP_FROM = 1 + 2
-// - OMP_MAP_BYCOPY = 128
-// CK14-DAG: [[TYPES:@.+]] = {{.+}}constant [2 x i32] [i32 3, i32 128]
+// - OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST = 35
+// - OMP_MAP_PRIVATE_VAL + OMP_MAP_IS_FIRST = 288
+// CK14-DAG: [[TYPES:@.+]] = {{.+}}constant [2 x i32] [i32 35, i32 288]
 
 class SSS {
 public:
@@ -743,15 +743,15 @@ void implicit_maps_class (int a){
 // CK15: [[ST:%.+]] = type { i32, double, i32* }
 // CK15: [[SIZES:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{24|16}}, i{{64|32}} 4]
 // Map types:
-// - OMP_MAP_TO | OMP_MAP_FROM = 1 + 2
-// - OMP_MAP_BYCOPY = 128
-// CK15: [[TYPES:@.+]] = {{.+}}constant [2 x i32] [i32 3, i32 128]
+// - OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST = 35
+// - OMP_MAP_PRIVATE_VAL + OMP_MAP_IS_FIRST = 288
+// CK15: [[TYPES:@.+]] = {{.+}}constant [2 x i32] [i32 35, i32 288]
 
 // CK15: [[SIZES2:@.+]] = {{.+}}constant [2 x i[[sz]]] [i{{64|32}} {{24|16}}, i{{64|32}} 4]
 // Map types:
-// - OMP_MAP_TO | OMP_MAP_FROM = 1 + 2
-// - OMP_MAP_BYCOPY = 128
-// CK15: [[TYPES2:@.+]] = {{.+}}constant [2 x i32] [i32 3, i32 128]
+// - OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST = 35
+// - OMP_MAP_PRIVATE_VAL + OMP_MAP_IS_FIRST = 288
+// CK15: [[TYPES2:@.+]] = {{.+}}constant [2 x i32] [i32 35, i32 288]
 
 template<int x>
 class SSST {
@@ -870,8 +870,8 @@ void implicit_maps_templated_class (int a){
 
 // CK16-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
 // Map types:
-// - OMP_MAP_BYCOPY = 128
-// CK16-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
+// - OMP_MAP_PRIVATE_VAL + OMP_MAP_IS_FIRST = 288
+// CK16-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 288]
 
 template<int y>
 int foo(int d) {
@@ -923,8 +923,8 @@ void implicit_maps_templated_function (int a){
 
 // CK17-DAG: [[ST:%.+]] = type { i32, double }
 // CK17-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} {{16|12}}]
-// Map types: OMP_MAP_TO + OMP_MAP_FROM = 2 + 1
-// CK17-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 3]
+// Map types: OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST = 35
+// CK17-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 35]
 
 class SSS {
 public:
@@ -971,8 +971,8 @@ void implicit_maps_struct (int a){
 
 // CK18-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
 // Map types:
-// - OMP_MAP_BYCOPY = 128
-// CK18-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
+// - OMP_MAP_PRIVATE_VAL + OMP_MAP_IS_FIRST = 288
+// CK18-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 288]
 
 template<typename T>
 int foo(T d) {
@@ -1022,122 +1022,122 @@ void implicit_maps_template_type_capture (int a){
 #ifdef CK19
 
 // CK19: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
-// CK19: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] zeroinitializer
+// CK19: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 32]
 
 // CK19: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400]
-// CK19: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 1]
+// CK19: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 33]
 
 // CK19: [[SIZE02:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 240]
-// CK19: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 2]
+// CK19: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 34]
 
 // CK19: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 240]
-// CK19: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK19: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK19: [[SIZE04:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400]
-// CK19: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] zeroinitializer
+// CK19: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 32]
 
 // CK19: [[SIZE05:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
-// CK19: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i32] [i32 1]
+// CK19: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i32] [i32 33]
 
-// CK19: [[MTYPE06:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK19: [[MTYPE06:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
-// CK19: [[MTYPE07:@.+]] = private {{.*}}constant [1 x i32] zeroinitializer
+// CK19: [[MTYPE07:@.+]] = private {{.*}}constant [1 x i32] [i32 32]
 
 // CK19: [[SIZE08:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
-// CK19: [[MTYPE08:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK19: [[MTYPE08:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK19: [[SIZE09:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{8|4}}]
-// CK19: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i32] [i32 2]
+// CK19: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i32] [i32 34]
 
 // CK19: [[SIZE10:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 240]
-// CK19: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK19: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK19: [[SIZE11:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 240]
-// CK19: [[MTYPE11:@.+]] = private {{.*}}constant [1 x i32] zeroinitializer
+// CK19: [[MTYPE11:@.+]] = private {{.*}}constant [1 x i32] [i32 32]
 
 // CK19: [[SIZE12:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
-// CK19: [[MTYPE12:@.+]] = private {{.*}}constant [1 x i32] [i32 1]
+// CK19: [[MTYPE12:@.+]] = private {{.*}}constant [1 x i32] [i32 33]
 
-// CK19: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i32] zeroinitializer
+// CK19: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i32] [i32 32]
 
-// CK19: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i32] [i32 1]
+// CK19: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i32] [i32 33]
 
 // CK19: [[SIZE15:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
-// CK19: [[MTYPE15:@.+]] = private {{.*}}constant [1 x i32] [i32 2]
+// CK19: [[MTYPE15:@.+]] = private {{.*}}constant [1 x i32] [i32 34]
 
-// CK19: [[MTYPE16:@.+]] = private {{.*}}constant [2 x i32] [i32 128, i32 1]
+// CK19: [[MTYPE16:@.+]] = private {{.*}}constant [2 x i32] [i32 288, i32 33]
 
 // CK19: [[SIZE17:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 240]
-// CK19: [[MTYPE17:@.+]] = private {{.*}}constant [2 x i32] [i32 128, i32 2]
+// CK19: [[MTYPE17:@.+]] = private {{.*}}constant [2 x i32] [i32 288, i32 34]
 
 // CK19: [[SIZE18:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 240]
-// CK19: [[MTYPE18:@.+]] = private {{.*}}constant [2 x i32] [i32 128, i32 3]
+// CK19: [[MTYPE18:@.+]] = private {{.*}}constant [2 x i32] [i32 288, i32 35]
 
-// CK19: [[MTYPE19:@.+]] = private {{.*}}constant [2 x i32] [i32 128, i32 0]
+// CK19: [[MTYPE19:@.+]] = private {{.*}}constant [2 x i32] [i32 288, i32 32]
 
 // CK19: [[SIZE20:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4]
-// CK19: [[MTYPE20:@.+]] = private {{.*}}constant [2 x i32] [i32 128, i32 1]
+// CK19: [[MTYPE20:@.+]] = private {{.*}}constant [2 x i32] [i32 288, i32 33]
 
-// CK19: [[MTYPE21:@.+]] = private {{.*}}constant [2 x i32] [i32 128, i32 3]
+// CK19: [[MTYPE21:@.+]] = private {{.*}}constant [2 x i32] [i32 288, i32 35]
 
 // CK19: [[SIZE22:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4]
-// CK19: [[MTYPE22:@.+]] = private {{.*}}constant [2 x i32] [i32 128, i32 3]
+// CK19: [[MTYPE22:@.+]] = private {{.*}}constant [2 x i32] [i32 288, i32 35]
 
 // CK19: [[SIZE23:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
-// CK19: [[MTYPE23:@.+]] = private {{.*}}constant [1 x i32] [i32 7]
+// CK19: [[MTYPE23:@.+]] = private {{.*}}constant [1 x i32] [i32 39]
 
 // CK19: [[SIZE24:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 480]
-// CK19: [[MTYPE24:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK19: [[MTYPE24:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK19: [[SIZE25:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 16]
-// CK19: [[MTYPE25:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK19: [[MTYPE25:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK19: [[SIZE26:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 24]
-// CK19: [[MTYPE26:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK19: [[MTYPE26:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK19: [[SIZE27:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
-// CK19: [[MTYPE27:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK19: [[MTYPE27:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK19: [[SIZE28:@.+]] = private {{.*}}constant [3 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] 16]
-// CK19: [[MTYPE28:@.+]] = private {{.*}}constant [3 x i32] [i32 3, i32 99, i32 99]
+// CK19: [[MTYPE28:@.+]] = private {{.*}}constant [3 x i32] [i32 35, i32 19, i32 19]
 
 // CK19: [[SIZE29:@.+]] = private {{.*}}constant [3 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] 4]
-// CK19: [[MTYPE29:@.+]] = private {{.*}}constant [3 x i32] [i32 3, i32 99, i32 99]
+// CK19: [[MTYPE29:@.+]] = private {{.*}}constant [3 x i32] [i32 35, i32 19, i32 19]
 
-// CK19: [[MTYPE30:@.+]] = private {{.*}}constant [4 x i32] [i32 128, i32 128, i32 128, i32 3]
+// CK19: [[MTYPE30:@.+]] = private {{.*}}constant [4 x i32] [i32 288, i32 288, i32 288, i32 35]
 
 // CK19: [[SIZE31:@.+]] = private {{.*}}constant [4 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] 40]
-// CK19: [[MTYPE31:@.+]] = private {{.*}}constant [4 x i32] [i32 128, i32 128, i32 128, i32 3]
+// CK19: [[MTYPE31:@.+]] = private {{.*}}constant [4 x i32] [i32 288, i32 288, i32 288, i32 35]
 
 // CK19: [[SIZE32:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 13728]
-// CK19: [[MTYPE32:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK19: [[MTYPE32:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK19: [[SIZE33:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 13728]
-// CK19: [[MTYPE33:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK19: [[MTYPE33:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK19: [[SIZE34:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 13728]
-// CK19: [[MTYPE34:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK19: [[MTYPE34:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
-// CK19: [[MTYPE35:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK19: [[MTYPE35:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK19: [[SIZE36:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 208]
-// CK19: [[MTYPE36:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK19: [[MTYPE36:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
-// CK19: [[MTYPE37:@.+]] = private {{.*}}constant [3 x i32] [i32 128, i32 128, i32 3]
+// CK19: [[MTYPE37:@.+]] = private {{.*}}constant [3 x i32] [i32 288, i32 288, i32 35]
 
-// CK19: [[MTYPE38:@.+]] = private {{.*}}constant [3 x i32] [i32 128, i32 128, i32 3]
+// CK19: [[MTYPE38:@.+]] = private {{.*}}constant [3 x i32] [i32 288, i32 288, i32 35]
 
-// CK19: [[MTYPE39:@.+]] = private {{.*}}constant [3 x i32] [i32 128, i32 128, i32 3]
+// CK19: [[MTYPE39:@.+]] = private {{.*}}constant [3 x i32] [i32 288, i32 288, i32 35]
 
-// CK19: [[MTYPE40:@.+]] = private {{.*}}constant [3 x i32] [i32 128, i32 128, i32 3]
+// CK19: [[MTYPE40:@.+]] = private {{.*}}constant [3 x i32] [i32 288, i32 288, i32 35]
 
 // CK19: [[SIZE41:@.+]] = private {{.*}}constant [3 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] 208]
-// CK19: [[MTYPE41:@.+]] = private {{.*}}constant [3 x i32] [i32 128, i32 128, i32 3]
+// CK19: [[MTYPE41:@.+]] = private {{.*}}constant [3 x i32] [i32 288, i32 288, i32 35]
 
 // CK19: [[SIZE42:@.+]] = private {{.*}}constant [3 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] 104]
-// CK19: [[MTYPE42:@.+]] = private {{.*}}constant [3 x i32] [i32 3, i32 99, i32 99]
+// CK19: [[MTYPE42:@.+]] = private {{.*}}constant [3 x i32] [i32 35, i32 19, i32 19]
 
-// CK19: [[MTYPE43:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK19: [[MTYPE43:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK19-LABEL: explicit_maps_single
 void explicit_maps_single (int ii){
@@ -2397,16 +2397,16 @@ void explicit_maps_single (int ii){
 #ifdef CK20
 
 // CK20: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
-// CK20: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 1]
+// CK20: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 33]
 
 // CK20: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20]
-// CK20: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 1]
+// CK20: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 33]
 
 // CK20: [[SIZE02:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
-// CK20: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 2]
+// CK20: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 34]
 
 // CK20: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 12]
-// CK20: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 2]
+// CK20: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 34]
 
 // CK20-LABEL: explicit_maps_references_and_function_args
 void explicit_maps_references_and_function_args (int a, float b, int (&c)[10], float *d){
@@ -2514,22 +2514,22 @@ void explicit_maps_references_and_function_args (int a, float b, int (&c)[10], f
 // CK21: [[ST:%.+]] = type { i32, i32, float* }
 
 // CK21: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
-// CK21: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK21: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK21: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 492]
-// CK21: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK21: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK21: [[SIZE02:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 500]
-// CK21: [[MTYPE02:@.+]] = private {{.*}}constant [2 x i32] [i32 2, i32 98]
+// CK21: [[MTYPE02:@.+]] = private {{.*}}constant [2 x i32] [i32 34, i32 18]
 
 // CK21: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 492]
-// CK21: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 2]
+// CK21: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 34]
 
 // CK21: [[SIZE04:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
-// CK21: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 2]
+// CK21: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 34]
 
 // CK21: [[SIZE05:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] 4, i[[Z]] 4]
-// CK21: [[MTYPE05:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 67]
+// CK21: [[MTYPE05:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 3]
 
 // CK21-LABEL: explicit_maps_template_args_and_members
 
@@ -2705,49 +2705,49 @@ int explicit_maps_template_args_and_members(int a){
 // CK22-DAG: [[STT:%.+]] = type { i32 }
 
 // CK22: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
-// CK22: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK22: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK22: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400]
-// CK22: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK22: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK22: [[SIZE02:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{8|4}}]
-// CK22: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK22: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK22: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 16]
-// CK22: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK22: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK22: [[SIZE04:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20]
-// CK22: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK22: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK22: [[SIZE05:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
-// CK22: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK22: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK22: [[SIZE06:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400]
-// CK22: [[MTYPE06:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK22: [[MTYPE06:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK22: [[SIZE07:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{8|4}}]
-// CK22: [[MTYPE07:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK22: [[MTYPE07:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK22: [[SIZE08:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 16]
-// CK22: [[MTYPE08:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK22: [[MTYPE08:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK22: [[SIZE09:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20]
-// CK22: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK22: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK22: [[SIZE10:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
-// CK22: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK22: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK22: [[SIZE11:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400]
-// CK22: [[MTYPE11:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK22: [[MTYPE11:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK22: [[SIZE12:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{8|4}}]
-// CK22: [[MTYPE12:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK22: [[MTYPE12:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK22: [[SIZE13:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 16]
-// CK22: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK22: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK22: [[SIZE14:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20]
-// CK22: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK22: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 int a;
 int c[100];
@@ -3025,22 +3025,22 @@ int explicit_maps_globals(void){
 #ifdef CK23
 
 // CK23: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
-// CK23: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK23: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK23: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
-// CK23: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK23: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK23: [[SIZE02:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400]
-// CK23: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK23: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK23: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{8|4}}]
-// CK23: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK23: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK23: [[SIZE04:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 16]
-// CK23: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK23: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK23: [[SIZE05:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 16]
-// CK23: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK23: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK23-LABEL: explicit_maps_inside_captured
 int explicit_maps_inside_captured(int a){
@@ -3215,76 +3215,76 @@ struct SC{
 };
 
 // CK24: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
-// CK24: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK24: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK24: [[SIZE02:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{56|48}}]
-// CK24: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK24: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK24: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
-// CK24: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK24: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK24: [[SIZE04:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20]
-// CK24: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK24: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK24: [[SIZE05:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{3560|2880}}]
-// CK24: [[MTYPE05:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99]
+// CK24: [[MTYPE05:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 19]
 
 // CK24: [[SIZE06:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
-// CK24: [[MTYPE06:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK24: [[MTYPE06:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK24: [[SIZE07:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4]
-// CK24: [[MTYPE07:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99]
+// CK24: [[MTYPE07:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 19]
 
 // CK24: [[SIZE08:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4]
-// CK24: [[MTYPE08:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99]
+// CK24: [[MTYPE08:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 19]
 
 // CK24: [[SIZE09:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4]
-// CK24: [[MTYPE09:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99]
+// CK24: [[MTYPE09:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 19]
 
 // CK24: [[SIZE10:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 8]
-// CK24: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK24: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK24: [[SIZE11:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}]
-// CK24: [[MTYPE11:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99]
+// CK24: [[MTYPE11:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 19]
 
 // CK24: [[SIZE12:@.+]] = private {{.*}}constant [4 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] 4]
-// CK24: [[MTYPE12:@.+]] = private {{.*}}constant [4 x i32] [i32 3, i32 99, i32 99, i32 99]
+// CK24: [[MTYPE12:@.+]] = private {{.*}}constant [4 x i32] [i32 35, i32 19, i32 19, i32 19]
 
 // CK24: [[SIZE13:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
-// CK24: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK24: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK24: [[SIZE14:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{56|48}}]
-// CK24: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK24: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK24: [[SIZE15:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
-// CK24: [[MTYPE15:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK24: [[MTYPE15:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK24: [[SIZE16:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20]
-// CK24: [[MTYPE16:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK24: [[MTYPE16:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK24: [[SIZE17:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{3560|2880}}]
-// CK24: [[MTYPE17:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99]
+// CK24: [[MTYPE17:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 19]
 
 // CK24: [[SIZE18:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
-// CK24: [[MTYPE18:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK24: [[MTYPE18:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK24: [[SIZE19:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4]
-// CK24: [[MTYPE19:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99]
+// CK24: [[MTYPE19:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 19]
 
 // CK24: [[SIZE20:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4]
-// CK24: [[MTYPE20:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99]
+// CK24: [[MTYPE20:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 19]
 
 // CK24: [[SIZE21:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4]
-// CK24: [[MTYPE21:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99]
+// CK24: [[MTYPE21:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 19]
 
 // CK24: [[SIZE22:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{8|4}}]
-// CK24: [[MTYPE22:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+// CK24: [[MTYPE22:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
 
 // CK24: [[SIZE23:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}]
-// CK24: [[MTYPE23:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99]
+// CK24: [[MTYPE23:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 19]
 
 // CK24: [[SIZE24:@.+]] = private {{.*}}constant [4 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] 4]
-// CK24: [[MTYPE24:@.+]] = private {{.*}}constant [4 x i32] [i32 3, i32 99, i32 99, i32 99]
+// CK24: [[MTYPE24:@.+]] = private {{.*}}constant [4 x i32] [i32 35, i32 19, i32 19, i32 19]
 
 // CK24-LABEL: explicit_maps_struct_fields
 int explicit_maps_struct_fields(int a){
@@ -3997,10 +3997,10 @@ int explicit_maps_struct_fields(int a){
 // CK25: [[CA01:%.+]] = type { i32* }
 
 // CK25: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
-// CK25: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 1]
+// CK25: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 33]
 
 // CK25: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
-// CK25: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 1]
+// CK25: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 33]
 
 // CK25-LABEL: explicit_maps_with_inner_lambda
 
@@ -4087,16 +4087,16 @@ int explicit_maps_with_inner_lambda(int a){
 // CK26: [[ST:%.+]] = type { i32, float*, i32, float* }
 
 // CK26: [[SIZE00:@.+]] = private {{.*}}constant [2 x i[[Z:64|32]]] [i[[Z:64|32]] {{32|16}}, i[[Z:64|32]] 4]
-// CK26: [[MTYPE00:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 3]
+// CK26: [[MTYPE00:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 35]
 
 // CK26: [[SIZE01:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{32|16}}, i[[Z]] 4]
-// CK26: [[MTYPE01:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 3]
+// CK26: [[MTYPE01:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 35]
 
 // CK26: [[SIZE02:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{32|16}}, i[[Z]] 4]
-// CK26: [[MTYPE02:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 3]
+// CK26: [[MTYPE02:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 35]
 
 // CK26: [[SIZE03:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{32|16}}, i[[Z]] 4]
-// CK26: [[MTYPE03:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 3]
+// CK26: [[MTYPE03:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 35]
 
 // CK26-LABEL: explicit_maps_with_private_class_members
 
@@ -4259,5 +4259,120 @@ int explicit_maps_with_private_class_members(){
   CC c(B);
   return c.foo();
 }
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK27 -verify -fopenmp -fomptargets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK27 --check-prefix CK27-64
+// RUN: %clang_cc1 -DCK27 -fopenmp -fomptargets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fomptargets=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 CK27 --check-prefix CK27-64
+// RUN: %clang_cc1 -DCK27 -verify -fopenmp -fomptargets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK27 --check-prefix CK27-32
+// RUN: %clang_cc1 -DCK27 -fopenmp -fomptargets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fomptargets=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 CK27 --check-prefix CK27-32
+#ifdef CK27
+
+// CK27: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] zeroinitializer
+// CK27: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 32]
+
+// CK27: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] zeroinitializer
+// CK27: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
+
+// CK27: [[SIZE02:@.+]] = private {{.*}}constant [1 x i[[Z]]] zeroinitializer
+// CK27: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
+
+// CK27: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] zeroinitializer
+// CK27: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
+
+// CK27-LABEL: zero_size_section_maps
+void zero_size_section_maps (int ii){
+
+  // Map of a pointer.
+  int *pa;
+
+  // Region 00
+  // CK27-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+  // 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: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK27-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK27-DAG: [[CBPVAL0]] = bitcast i32* [[VAR0:%.+]] to i8*
+  // CK27-DAG: [[CPVAL0]] = bitcast i32* [[VAR0]] to i8*
+
+  // CK27: call void [[CALL00:@.+]](i32* {{[^,]+}})
+  #pragma omp target
+  {
+    pa[50]++;
+  }
+
+  // Region 01
+  // CK27-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
+  // 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: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK27-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK27-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8*
+  // CK27-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
+  // 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(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}})
+  // 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: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK27-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK27-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8*
+  // CK27-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
+  // 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(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE03]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}})
+  // 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: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK27-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK27-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8*
+  // CK27-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
+  // 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]++;
+  }
+}
+
+// CK27: define {{.+}}[[CALL00]]
+// CK27: define {{.+}}[[CALL01]]
+// CK27: define {{.+}}[[CALL02]]
+// CK27: define {{.+}}[[CALL03]]
+
 #endif
 #endif