Index: lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- lib/CodeGen/CGOpenMPRuntime.cpp
+++ lib/CodeGen/CGOpenMPRuntime.cpp
@@ -4858,8 +4858,6 @@
   /// \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.
@@ -4867,19 +4865,19 @@
     /// \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;
@@ -4936,14 +4934,19 @@
 
   /// \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;
@@ -4957,17 +4960,14 @@
     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;
@@ -5219,13 +5219,13 @@
 
         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)
@@ -5532,7 +5532,8 @@
       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.
@@ -5551,10 +5552,10 @@
           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");
@@ -5573,11 +5574,17 @@
             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);
@@ -5589,13 +5596,15 @@
           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.
Index: test/OpenMP/target_codegen.cpp
===================================================================
--- test/OpenMP/target_codegen.cpp
+++ 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
Index: test/OpenMP/target_codegen_registration.cpp
===================================================================
--- test/OpenMP/target_codegen_registration.cpp
+++ 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
Index: test/OpenMP/target_data_codegen.cpp
===================================================================
--- test/OpenMP/target_data_codegen.cpp
+++ test/OpenMP/target_data_codegen.cpp
@@ -22,15 +22,15 @@
 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 @@
 };
 
 // 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){
Index: test/OpenMP/target_enter_data_codegen.cpp
===================================================================
--- test/OpenMP/target_enter_data_codegen.cpp
+++ test/OpenMP/target_enter_data_codegen.cpp
@@ -22,15 +22,15 @@
 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 @@
 };
 
 // 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){
Index: test/OpenMP/target_exit_data_codegen.cpp
===================================================================
--- test/OpenMP/target_exit_data_codegen.cpp
+++ test/OpenMP/target_exit_data_codegen.cpp
@@ -22,15 +22,15 @@
 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 @@
 };
 
 // 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){
Index: test/OpenMP/target_firstprivate_codegen.cpp
===================================================================
--- test/OpenMP/target_firstprivate_codegen.cpp
+++ test/OpenMP/target_firstprivate_codegen.cpp
@@ -33,16 +33,15 @@
 // 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:@.+]](
Index: test/OpenMP/target_map_codegen.cpp
===================================================================
--- test/OpenMP/target_map_codegen.cpp
+++ 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 @@
 #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 @@
 #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 @@
 #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 @@
 #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 @@
 #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 @@
 // 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 @@
 #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 @@
 #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 @@
 // 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 @@
 #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 @@
 // 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 @@
 
 // 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 @@
 // 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 @@
 // 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 @@
 
 // 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 @@
 
 // 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 @@
 
 // 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 @@
 #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 @@
 #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 @@
 // 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 @@
 // 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 @@
 #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 @@
 };
 
 // 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 @@
 // 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 @@
 // 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
 
@@ -4260,4 +4260,119 @@
   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