[OpenMP] Map clause codegeneration.

[sfantao]

Summary

Implement codegen for the map clause. All the new list items in 4.5 specification are supported.

Fix bug in the generation of array sections that was exposed by some of the map clause tests: for pointer types the offsets have to be calculated from the pointee not the pointer.

[sfantao]

Ping!

[ABataev]

lib/CodeGen/CGExpr.cpp

2937–2943

This must be fixed already, please reabse

lib/CodeGen/CGOpenMPRuntime.cpp

3915–3946

I don't think this is correct. It won't work for 2(and more)-dimensional array sections. Instead calculate size as 'last_item_address-first_item_address+1'.

4126–4128

In a few days you won't need it at all. All member references are implicitly captured into special declaration and you should not worry about them at all

[sfantao]

Rebase and add more regression tests for multidimensional arrays and arrays of pointers.

[sfantao]

Hi Alexey,

Thanks for the review!

lib/CodeGen/CGExpr.cpp

2937–2943

Thanks for the fix! I've just posted a rebased patch.

lib/CodeGen/CGOpenMPRuntime.cpp

3915–3946

I don't see why it should not work. There was only a small issue related with array decays to pointer that I fixed in the new diff. The size of the section always refer to the right most expression, even for multidimensional arrays. I added regression tests for that.

Other reason I avoided computing sizes based on addresses is that they do not get folded into constants. In code generation we have optimized behaviors for when all the sizes are constants so that a constant array is implemented instead of filling (at runtime) an array with constants to be passed to the runtime library.

Let me know if you disagree.

4126–4128

Great! Thanks for working on that! I'll then see how that will interfere with the map clause code generation.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

3915–3946

What if the base is OMPArraySectionExpr? Will it work in this case?

[sfantao]

Hi Alexey,

lib/CodeGen/CGOpenMPRuntime.cpp

3915–3946

Yes, it will work. There are a few regression tests for that in this patch already.

[sfantao]

test/OpenMP/target_map_codegen.cpp

1705

Regression tests for array sections start here.

[sfantao]

Rebase.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

3915–3946

Did not find any with array sections as base. Could please point a test where I can find something like 'marr[:][0:2][0:2]' as a mapped expression?
Also, it will produce some recalculation for VLAs. Add a test for it, please.

[sfantao]

Add regression tests for invalid array sections and multidimensional VLA.

[sfantao]

Hi Alexey,

Thanks for the feedback!

lib/CodeGen/CGOpenMPRuntime.cpp

3915–3946

Did not find any with array sections as base. Could please point a test where I can find something like 'marr[:][0:2][0:2]' as a mapped expression?

Oh, a map that uses an array expression that is not in the rightmost subexpression is illegal because it would correspond to non-contiguous storage. Sema is already taking care of that. I added your specific example to the Sema tests anyway.

Also, it will produce some recalculation for VLAs. Add a test for it, please.

Added tests for multidimensional VLAs.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

3915–3946

'marr[:][:][:2]' produces contiguous storage, no? Is this allowed? And if it is allowed, how it will be handled? Besides, if 'marr' is a pointer, even 'marr[:2][:3][:4]' may result in contiguous space.

[sfantao]

Hi Alexey,

lib/CodeGen/CGOpenMPRuntime.cpp

3915–3946

'marr[:][:][:2]' produces contiguous storage, no? Is this allowed? And if it is allowed, how it will be handled?

That is not contiguous so it is not allowed. marr[:2][:][:] is contiguous storage but can be expressed as marr[:2] only. Do you think I should create a patch to allow the latter case in Sema? Should it be part of this patch?

Besides, if 'marr' is a pointer, even 'marr[:2][:3][:4]' may result in contiguous space.

This is illegal. This requires multiple memory allocations that are not contiguous on the host.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

3915–3946

That is not contiguous so it is not allowed. marr[:2][:][:] is contiguous storage but can be expressed as marr[:2] only. Do you think I should create a patch to allow the latter case in Sema? Should it be part of this patch?

Yes, I think this must be supported. Also we should support marr[:2][0:size2][0:size3] forms, if they produce contiguous space. No, it should be a separate patch.

[sfantao]

Add support for valid array sections with multiple dimensions in map clause.

lib/CodeGen/CGOpenMPRuntime.cpp

3915–3946

Done, posted a new patch for the SEMA support and updated this one with regression tests for multidimensional array sections and few changes in map elements scanning to handle array sections correctly.

[sfantao]

Rebase and use getBaseOriginalType.

[sfantao]

• Rebase.
• Add regression test for non-contiguous storage map.

[sfantao]

Rebase.

[sfantao]

Ping!

[ABataev]

Samuel, I have a patch that improves handling of map clauses with data members. Unfortunately, I cannot commit it because I can't create tests for it. Please, apply this patch to your map clause codegen changes. It uses common technique for handling non-static data members:

diff --git a/include/clang/AST/OpenMPClause.h b/include/clang/AST/OpenMPClause.h
index 48a03b1..f51fadb 100644
--- a/include/clang/AST/OpenMPClause.h
+++ b/include/clang/AST/OpenMPClause.h
@@ -2784,6 +2784,7 @@ public:
 /// with the variables 'a' and 'b'.
 ///
 class OMPMapClause final : public OMPVarListClause<OMPMapClause>,
+                           public OMPClauseWithPreInit,
                            private llvm::TrailingObjects<OMPMapClause, Expr *> {
   friend TrailingObjects;
   friend OMPVarListClause;
@@ -2838,8 +2839,9 @@ class OMPMapClause final : public OMPVarListClause<OMPMapClause>,
                         unsigned N)
       : OMPVarListClause<OMPMapClause>(OMPC_map, StartLoc, LParenLoc, EndLoc,
                                        N),
-        MapTypeModifier(MapTypeModifier), MapType(MapType),
-        MapTypeIsImplicit(MapTypeIsImplicit), MapLoc(MapLoc) {}
+        OMPClauseWithPreInit(this), MapTypeModifier(MapTypeModifier),
+        MapType(MapType), MapTypeIsImplicit(MapTypeIsImplicit), MapLoc(MapLoc) {
+  }

   /// \brief Build an empty clause.
   ///
@@ -2848,8 +2850,8 @@ class OMPMapClause final : public OMPVarListClause<OMPMapClause>,
   explicit OMPMapClause(unsigned N)
       : OMPVarListClause<OMPMapClause>(OMPC_map, SourceLocation(),
                                        SourceLocation(), SourceLocation(), N),
-        MapTypeModifier(OMPC_MAP_unknown), MapType(OMPC_MAP_unknown),
-        MapTypeIsImplicit(false), MapLoc() {}
+        OMPClauseWithPreInit(this), MapTypeModifier(OMPC_MAP_unknown),
+        MapType(OMPC_MAP_unknown), MapTypeIsImplicit(false), MapLoc() {}

 public:
   /// \brief Creates clause with a list of variables \a VL.
@@ -2862,13 +2864,15 @@ public:
   /// \param Type Map type.
   /// \param TypeIsImplicit Map type is inferred implicitly.
   /// \param TypeLoc Location of the map type.
+  /// \param PreInit Statement that must be executed before entering the OpenMP
+  /// region with this clause.
   ///
   static OMPMapClause *Create(const ASTContext &C, SourceLocation StartLoc,
                               SourceLocation LParenLoc, SourceLocation EndLoc,
                               ArrayRef<Expr *> VL,
                               OpenMPMapClauseKind TypeModifier,
                               OpenMPMapClauseKind Type, bool TypeIsImplicit,
-                              SourceLocation TypeLoc);
+                              SourceLocation TypeLoc, Stmt *PreInit);
   /// \brief Creates an empty clause with the place for \a N variables.
   ///
   /// \param C AST context.
diff --git a/include/clang/AST/RecursiveASTVisitor.h b/include/clang/AST/RecursiveASTVisitor.h
index 6d86727..7835a17 100644
--- a/include/clang/AST/RecursiveASTVisitor.h
+++ b/include/clang/AST/RecursiveASTVisitor.h
@@ -2816,6 +2816,7 @@ bool RecursiveASTVisitor<Derived>::VisitOMPDeviceClause(OMPDeviceClause *C) {
 template <typename Derived>
 bool RecursiveASTVisitor<Derived>::VisitOMPMapClause(OMPMapClause *C) {
   TRY_TO(VisitOMPClauseList(C));
+  TRY_TO(VisitOMPClauseWithPreInit(C));
   return true;
 }

diff --git a/lib/AST/OpenMPClause.cpp b/lib/AST/OpenMPClause.cpp
index 8843ded..8cc6b86 100644
--- a/lib/AST/OpenMPClause.cpp
+++ b/lib/AST/OpenMPClause.cpp
@@ -48,6 +48,8 @@ const OMPClauseWithPreInit *OMPClauseWithPreInit::get(const OMPClause *C) {
     return static_cast<const OMPReductionClause *>(C);
   case OMPC_linear:
     return static_cast<const OMPLinearClause *>(C);
+  case OMPC_map:
+    return static_cast<const OMPMapClause *>(C);
   case OMPC_default:
   case OMPC_proc_bind:
   case OMPC_if:
@@ -76,7 +78,6 @@ const OMPClauseWithPreInit *OMPClauseWithPreInit::get(const OMPClause *C) {
   case OMPC_device:
   case OMPC_threads:
   case OMPC_simd:
-  case OMPC_map:
   case OMPC_num_teams:
   case OMPC_thread_limit:
   case OMPC_priority:
@@ -533,8 +534,8 @@ OMPMapClause *OMPMapClause::Create(const ASTContext &C, SourceLocation StartLoc,
                                    SourceLocation EndLoc, ArrayRef<Expr *> VL,
                                    OpenMPMapClauseKind TypeModifier,
                                    OpenMPMapClauseKind Type,
-                                   bool TypeIsImplicit,
-                                   SourceLocation TypeLoc) {
+                                   bool TypeIsImplicit, SourceLocation TypeLoc,
+                                   Stmt *PreInit) {
   void *Mem = C.Allocate(totalSizeToAlloc<Expr *>(VL.size()));
   OMPMapClause *Clause =
       new (Mem) OMPMapClause(TypeModifier, Type, TypeIsImplicit, TypeLoc,
@@ -543,6 +544,7 @@ OMPMapClause *OMPMapClause::Create(const ASTContext &C, SourceLocation StartLoc,
   Clause->setMapTypeModifier(TypeModifier);
   Clause->setMapType(Type);
   Clause->setMapLoc(TypeLoc);
+  Clause->setPreInitStmt(PreInit);
   return Clause;
 }

diff --git a/lib/AST/StmtProfile.cpp b/lib/AST/StmtProfile.cpp
index 327b4c0..68fb535 100644
--- a/lib/AST/StmtProfile.cpp
+++ b/lib/AST/StmtProfile.cpp
@@ -471,6 +471,7 @@ void OMPClauseProfiler::VisitOMPDeviceClause(const OMPDeviceClause *C) {
 }
 void OMPClauseProfiler::VisitOMPMapClause(const OMPMapClause *C) {
   VisitOMPClauseList(C);
+  VistOMPClauseWithPreInit(C);
 }
 void OMPClauseProfiler::VisitOMPNumTeamsClause(const OMPNumTeamsClause *C) {
   Profiler->VisitStmt(C->getNumTeams());
diff --git a/lib/Sema/SemaOpenMP.cpp b/lib/Sema/SemaOpenMP.cpp
index 826f069..7d3e3cf 100644
--- a/lib/Sema/SemaOpenMP.cpp
+++ b/lib/Sema/SemaOpenMP.cpp
@@ -9776,16 +9776,11 @@ Sema::ActOnOpenMPMapClause(OpenMPMapClauseKind MapTypeModifier,
                            ArrayRef<Expr *> VarList, SourceLocation StartLoc,
                            SourceLocation LParenLoc, SourceLocation EndLoc) {
   SmallVector<Expr *, 4> Vars;
+  SmallVector<Decl *, 4> ExprCaptures;

   for (auto &RE : VarList) {
     assert(RE && "Null expr in omp map");
-    if (isa<DependentScopeDeclRefExpr>(RE)) {
-      // It will be analyzed later.
-      Vars.push_back(RE);
-      continue;
-    }
     SourceLocation ELoc = RE->getExprLoc();
-
     auto *VE = RE->IgnoreParenLValueCasts();

     if (VE->isValueDependent() || VE->isTypeDependent() ||
@@ -9797,9 +9792,9 @@ Sema::ActOnOpenMPMapClause(OpenMPMapClauseKind MapTypeModifier,
       continue;
     }

-    auto *SimpleExpr = RE->IgnoreParenCasts();
+    auto *SimpleExpr = VE->IgnoreParenCasts();

-    if (!RE->IgnoreParenImpCasts()->isLValue()) {
+    if (!VE->isLValue()) {
       Diag(ELoc, diag::err_omp_expected_named_var_member_or_array_expression)
           << RE->getSourceRange();
       continue;
@@ -9897,15 +9892,43 @@ Sema::ActOnOpenMPMapClause(OpenMPMapClauseKind MapTypeModifier,
       continue;
     }

-    Vars.push_back(RE);
-    DSAStack->addExprToVarMapInfo(D, RE);
+    DeclRefExpr *Ref = nullptr;
+    if (!VD)
+      Ref = buildCapture(*this, D, SimpleRefExpr, /*WithInit=*/true);
+    DeclRefExpr *Ref = nullptr;
+    Expr *VarsExpr = RE->IgnoreParens();
+    if (!VD) {
+      if (BE != SimpleExpr) {
+        TransformExprToCaptures RebuildToCapture(*this, D);
+        VarsExpr =
+            RebuildToCapture.TransformExpr(RefExpr->IgnoreParens()).get();
+        Ref = RebuildToCapture.getCapturedExpr();
+      } else {
+        VarsExpr = Ref =
+            buildCapture(*this, D, SimpleRefExpr, /*WithInit=*/true);
+      }
+      if (!IsOpenMPCapturedDecl(D))
+        ExprCaptures.push_back(Ref->getDecl());
+    }
+    Vars.push_back(VarsExpr);
+    DSAStack->addExprToVarMapInfo(D, RE->IgnoreParens());
   }

+  if (Vars.empty())
+    return nullptr;
+
+  Stmt *PreInit = nullptr;
+  if (!ExprCaptures.empty()) {
+    PreInit = new (Context)
+        DeclStmt(DeclGroupRef::Create(Context, ExprCaptures.begin(),
+                                      ExprCaptures.size()),
+                 SourceLocation(), SourceLocation());
+  }
   // We need to produce a map clause even if we don't have variables so that
   // other diagnostics related with non-existing map clauses are accurate.
   return OMPMapClause::Create(Context, StartLoc, LParenLoc, EndLoc, Vars,
                               MapTypeModifier, MapType, IsMapTypeImplicit,
-                              MapLoc);
+                              MapLoc, PreInit);
 }

 QualType Sema::ActOnOpenMPDeclareReductionType(SourceLocation TyLoc,
diff --git a/lib/Serialization/ASTReaderStmt.cpp b/lib/Serialization/ASTReaderStmt.cpp
index ab44c96..941caf7 100644
--- a/lib/Serialization/ASTReaderStmt.cpp
+++ b/lib/Serialization/ASTReaderStmt.cpp
@@ -2214,6 +2214,7 @@ void OMPClauseReader::VisitOMPDeviceClause(OMPDeviceClause *C) {
 }

 void OMPClauseReader::VisitOMPMapClause(OMPMapClause *C) {
+  VisitOMPClauseWithPreInit(C);
   C->setLParenLoc(Reader->ReadSourceLocation(Record, Idx));
   C->setMapTypeModifier(
      static_cast<OpenMPMapClauseKind>(Record[Idx++]));
diff --git a/lib/Serialization/ASTWriterStmt.cpp b/lib/Serialization/ASTWriterStmt.cpp
index ea3b804..6d3dca3 100644
--- a/lib/Serialization/ASTWriterStmt.cpp
+++ b/lib/Serialization/ASTWriterStmt.cpp
@@ -2013,6 +2013,7 @@ void OMPClauseWriter::VisitOMPDeviceClause(OMPDeviceClause *C) {

 void OMPClauseWriter::VisitOMPMapClause(OMPMapClause *C) {
   Record.push_back(C->varlist_size());
+  VisitOMPClauseWithPreInit(C);
   Writer->Writer.AddSourceLocation(C->getLParenLoc(), Record);
   Record.push_back(C->getMapTypeModifier());
   Record.push_back(C->getMapType());
diff --git a/tools/libclang/CIndex.cpp b/tools/libclang/CIndex.cpp
index 5dca6f9..05c0380 100644
--- a/tools/libclang/CIndex.cpp
+++ b/tools/libclang/CIndex.cpp
@@ -2247,6 +2247,7 @@ void OMPClauseEnqueue::VisitOMPDependClause(const OMPDependClause *C) {
 }
 void OMPClauseEnqueue::VisitOMPMapClause(const OMPMapClause *C) {
   VisitOMPClauseList(C);
+  VisitOMPClauseWithPreInit(C);
 }
 void OMPClauseEnqueue::VisitOMPDistScheduleClause(
     const OMPDistScheduleClause *C) {

[sfantao]

Hi Alexey,

I am afraid I don't understand what you are trying to accomplish with this. I was unable to compile your patch as is, I think what you intended is something like this:

// If we are mapping a field of 'this' we attempt to generate a
// OMPCapturedExprDecl for it. The logic above, already ensures that if we
// map a member expression its base is 'this'.
Expr *VarsExpr = RE->IgnoreParens();
if (isa<MemberExpr>(BE)) {
  DeclRefExpr *Ref;
  VarsExpr = Ref = buildCapture(*this, D, BE, /*WithInit=*/true);
  if (!IsOpenMPCapturedDecl(D))
    ExprCaptures.push_back(Ref->getDecl());
}

However, I am not sure how this helps in this case. The map clause will have expressions that refer to the dummy declaration, but the region captured arguments are unchanged, so we still get 'this' and the content of the map clause has to relate to that.

Say we have some user code like this:

struct S {
  int A;
  foo() { ++ A; }
  bar() {
    #pragma target map(A)
    {
      ++A;
      foo()
    }
}

The kernel argument is going to be this regardless the map clause. The map clause only makes sure that a section of this is copied and not everything, so I rather not separate the field from the rest (for the same reason we do not separate the array declaration from a subscript or section). this still has to be passed to foo().

Is it your goal to pass arguments to the outlined region individually? I don't think your code was trying to do that, am I wrong? Even if we manage to change the captures so that fields are passed individually, this would have to be rematerialized inside bar().

Am I missing something?

Thanks!
Samuel

[ABataev]

Samuel, this will be captured. But foo() will use original A. This cannot be changed.
But Inside the region all reference to this->A will be replaced by the CapturedDecl A. It will allow proper codegen for inner captured contexts.
For example, such code won't be handled correctly without my changes:
#pragma omp target map(to: this->A)
[]()->void {++A}();

In this code you will work with original A, not the one captured in the target region.

I reworked the patch a little bit, here is improved version:

diff --git a/include/clang/AST/OpenMPClause.h b/include/clang/AST/OpenMPClause.h
index 48a03b1..f51fadb 100644
--- a/include/clang/AST/OpenMPClause.h
+++ b/include/clang/AST/OpenMPClause.h
@@ -2784,6 +2784,7 @@ public:
 /// with the variables 'a' and 'b'.
 ///
 class OMPMapClause final : public OMPVarListClause<OMPMapClause>,
+                           public OMPClauseWithPreInit,
                            private llvm::TrailingObjects<OMPMapClause, Expr *> {
   friend TrailingObjects;
   friend OMPVarListClause;
@@ -2838,8 +2839,9 @@ class OMPMapClause final : public OMPVarListClause<OMPMapClause>,
                         unsigned N)
       : OMPVarListClause<OMPMapClause>(OMPC_map, StartLoc, LParenLoc, EndLoc,
                                        N),
-        MapTypeModifier(MapTypeModifier), MapType(MapType),
-        MapTypeIsImplicit(MapTypeIsImplicit), MapLoc(MapLoc) {}
+        OMPClauseWithPreInit(this), MapTypeModifier(MapTypeModifier),
+        MapType(MapType), MapTypeIsImplicit(MapTypeIsImplicit), MapLoc(MapLoc) {
+  }

   /// \brief Build an empty clause.
   ///
@@ -2848,8 +2850,8 @@ class OMPMapClause final : public OMPVarListClause<OMPMapClause>,
   explicit OMPMapClause(unsigned N)
       : OMPVarListClause<OMPMapClause>(OMPC_map, SourceLocation(),
                                        SourceLocation(), SourceLocation(), N),
-        MapTypeModifier(OMPC_MAP_unknown), MapType(OMPC_MAP_unknown),
-        MapTypeIsImplicit(false), MapLoc() {}
+        OMPClauseWithPreInit(this), MapTypeModifier(OMPC_MAP_unknown),
+        MapType(OMPC_MAP_unknown), MapTypeIsImplicit(false), MapLoc() {}

 public:
   /// \brief Creates clause with a list of variables \a VL.
@@ -2862,13 +2864,15 @@ public:
   /// \param Type Map type.
   /// \param TypeIsImplicit Map type is inferred implicitly.
   /// \param TypeLoc Location of the map type.
+  /// \param PreInit Statement that must be executed before entering the OpenMP
+  /// region with this clause.
   ///
   static OMPMapClause *Create(const ASTContext &C, SourceLocation StartLoc,
                               SourceLocation LParenLoc, SourceLocation EndLoc,
                               ArrayRef<Expr *> VL,
                               OpenMPMapClauseKind TypeModifier,
                               OpenMPMapClauseKind Type, bool TypeIsImplicit,
-                              SourceLocation TypeLoc);
+                              SourceLocation TypeLoc, Stmt *PreInit);
   /// \brief Creates an empty clause with the place for \a N variables.
   ///
   /// \param C AST context.
diff --git a/include/clang/AST/RecursiveASTVisitor.h b/include/clang/AST/RecursiveASTVisitor.h
index 6d86727..7835a17 100644
--- a/include/clang/AST/RecursiveASTVisitor.h
+++ b/include/clang/AST/RecursiveASTVisitor.h
@@ -2816,6 +2816,7 @@ bool RecursiveASTVisitor<Derived>::VisitOMPDeviceClause(OMPDeviceClause *C) {
 template <typename Derived>
 bool RecursiveASTVisitor<Derived>::VisitOMPMapClause(OMPMapClause *C) {
   TRY_TO(VisitOMPClauseList(C));
+  TRY_TO(VisitOMPClauseWithPreInit(C));
   return true;
 }

diff --git a/lib/AST/OpenMPClause.cpp b/lib/AST/OpenMPClause.cpp
index 8843ded..8cc6b86 100644
--- a/lib/AST/OpenMPClause.cpp
+++ b/lib/AST/OpenMPClause.cpp
@@ -48,6 +48,8 @@ const OMPClauseWithPreInit *OMPClauseWithPreInit::get(const OMPClause *C) {
     return static_cast<const OMPReductionClause *>(C);
   case OMPC_linear:
     return static_cast<const OMPLinearClause *>(C);
+  case OMPC_map:
+    return static_cast<const OMPMapClause *>(C);
   case OMPC_default:
   case OMPC_proc_bind:
   case OMPC_if:
@@ -76,7 +78,6 @@ const OMPClauseWithPreInit *OMPClauseWithPreInit::get(const OMPClause *C) {
   case OMPC_device:
   case OMPC_threads:
   case OMPC_simd:
-  case OMPC_map:
   case OMPC_num_teams:
   case OMPC_thread_limit:
   case OMPC_priority:
@@ -533,8 +534,8 @@ OMPMapClause *OMPMapClause::Create(const ASTContext &C, SourceLocation StartLoc,
                                    SourceLocation EndLoc, ArrayRef<Expr *> VL,
                                    OpenMPMapClauseKind TypeModifier,
                                    OpenMPMapClauseKind Type,
-                                   bool TypeIsImplicit,
-                                   SourceLocation TypeLoc) {
+                                   bool TypeIsImplicit, SourceLocation TypeLoc,
+                                   Stmt *PreInit) {
   void *Mem = C.Allocate(totalSizeToAlloc<Expr *>(VL.size()));
   OMPMapClause *Clause =
       new (Mem) OMPMapClause(TypeModifier, Type, TypeIsImplicit, TypeLoc,
@@ -543,6 +544,7 @@ OMPMapClause *OMPMapClause::Create(const ASTContext &C, SourceLocation StartLoc,
   Clause->setMapTypeModifier(TypeModifier);
   Clause->setMapType(Type);
   Clause->setMapLoc(TypeLoc);
+  Clause->setPreInitStmt(PreInit);
   return Clause;
 }

diff --git a/lib/AST/StmtProfile.cpp b/lib/AST/StmtProfile.cpp
index 327b4c0..68fb535 100644
--- a/lib/AST/StmtProfile.cpp
+++ b/lib/AST/StmtProfile.cpp
@@ -471,6 +471,7 @@ void OMPClauseProfiler::VisitOMPDeviceClause(const OMPDeviceClause *C) {
 }
 void OMPClauseProfiler::VisitOMPMapClause(const OMPMapClause *C) {
   VisitOMPClauseList(C);
+  VistOMPClauseWithPreInit(C);
 }
 void OMPClauseProfiler::VisitOMPNumTeamsClause(const OMPNumTeamsClause *C) {
   Profiler->VisitStmt(C->getNumTeams());
diff --git a/lib/Sema/SemaOpenMP.cpp b/lib/Sema/SemaOpenMP.cpp
index 826f069..7d3e3cf 100644
--- a/lib/Sema/SemaOpenMP.cpp
+++ b/lib/Sema/SemaOpenMP.cpp
@@ -9776,16 +9776,11 @@ Sema::ActOnOpenMPMapClause(OpenMPMapClauseKind MapTypeModifier,
                            ArrayRef<Expr *> VarList, SourceLocation StartLoc,
                            SourceLocation LParenLoc, SourceLocation EndLoc) {
   SmallVector<Expr *, 4> Vars;
+  SmallVector<Decl *, 4> ExprCaptures;

   for (auto &RE : VarList) {
     assert(RE && "Null expr in omp map");
-    if (isa<DependentScopeDeclRefExpr>(RE)) {
-      // It will be analyzed later.
-      Vars.push_back(RE);
-      continue;
-    }
     SourceLocation ELoc = RE->getExprLoc();
-
     auto *VE = RE->IgnoreParenLValueCasts();

     if (VE->isValueDependent() || VE->isTypeDependent() ||
@@ -9797,9 +9792,9 @@ Sema::ActOnOpenMPMapClause(OpenMPMapClauseKind MapTypeModifier,
       continue;
     }

-    auto *SimpleExpr = RE->IgnoreParenCasts();
+    auto *SimpleExpr = VE->IgnoreParenCasts();

-    if (!RE->IgnoreParenImpCasts()->isLValue()) {
+    if (!VE->isLValue()) {
       Diag(ELoc, diag::err_omp_expected_named_var_member_or_array_expression)
           << RE->getSourceRange();
       continue;
@@ -9897,15 +9892,43 @@ Sema::ActOnOpenMPMapClause(OpenMPMapClauseKind MapTypeModifier,
       continue;
     }

-    Vars.push_back(RE);
-    DSAStack->addExprToVarMapInfo(D, RE);
+    Expr *VarsExpr = RE->IgnoreParens();
+    if (!VD) {
+      DeclRefExpr *Ref;
+      if (BE != SimpleExpr) {
+        TransformExprToCaptures RebuildToCapture(*this, D);
+        VarsExpr =
+            RebuildToCapture.TransformExpr(RefExpr->IgnoreParens()).get();
+        Ref = RebuildToCapture.getCapturedExpr();
+      } else {
+        VarsExpr = Ref =
+            buildCapture(*this, D, SimpleRefExpr, /*WithInit=*/true);
+      }
+      if (!IsOpenMPCapturedDecl(D))
+        ExprCaptures.push_back(Ref->getDecl());
+    }
+    Vars.push_back(VarsExpr);
+    DSAStack->addExprToVarMapInfo(D, RE->IgnoreParens());
   }

+  if (Vars.empty())
+    return nullptr;
+
+  Stmt *PreInit = nullptr;
+  if (!ExprCaptures.empty()) {
+    PreInit = new (Context)
+        DeclStmt(DeclGroupRef::Create(Context, ExprCaptures.begin(),
+                                      ExprCaptures.size()),
+                 SourceLocation(), SourceLocation());
+  }
   // We need to produce a map clause even if we don't have variables so that
   // other diagnostics related with non-existing map clauses are accurate.
   return OMPMapClause::Create(Context, StartLoc, LParenLoc, EndLoc, Vars,
                               MapTypeModifier, MapType, IsMapTypeImplicit,
-                              MapLoc);
+                              MapLoc, PreInit);
 }

 QualType Sema::ActOnOpenMPDeclareReductionType(SourceLocation TyLoc,
diff --git a/lib/Serialization/ASTReaderStmt.cpp b/lib/Serialization/ASTReaderStmt.cpp
index ab44c96..941caf7 100644
--- a/lib/Serialization/ASTReaderStmt.cpp
+++ b/lib/Serialization/ASTReaderStmt.cpp
@@ -2214,6 +2214,7 @@ void OMPClauseReader::VisitOMPDeviceClause(OMPDeviceClause *C) {
 }

 void OMPClauseReader::VisitOMPMapClause(OMPMapClause *C) {
+  VisitOMPClauseWithPreInit(C);
   C->setLParenLoc(Reader->ReadSourceLocation(Record, Idx));
   C->setMapTypeModifier(
      static_cast<OpenMPMapClauseKind>(Record[Idx++]));
diff --git a/lib/Serialization/ASTWriterStmt.cpp b/lib/Serialization/ASTWriterStmt.cpp
index ea3b804..6d3dca3 100644
--- a/lib/Serialization/ASTWriterStmt.cpp
+++ b/lib/Serialization/ASTWriterStmt.cpp
@@ -2013,6 +2013,7 @@ void OMPClauseWriter::VisitOMPDeviceClause(OMPDeviceClause *C) {

 void OMPClauseWriter::VisitOMPMapClause(OMPMapClause *C) {
   Record.push_back(C->varlist_size());
+  VisitOMPClauseWithPreInit(C);
   Writer->Writer.AddSourceLocation(C->getLParenLoc(), Record);
   Record.push_back(C->getMapTypeModifier());
   Record.push_back(C->getMapType());
diff --git a/tools/libclang/CIndex.cpp b/tools/libclang/CIndex.cpp
index 5dca6f9..05c0380 100644
--- a/tools/libclang/CIndex.cpp
+++ b/tools/libclang/CIndex.cpp
@@ -2247,6 +2247,7 @@ void OMPClauseEnqueue::VisitOMPDependClause(const OMPDependClause *C) {
 }
 void OMPClauseEnqueue::VisitOMPMapClause(const OMPMapClause *C) {
   VisitOMPClauseList(C);
+  VisitOMPClauseWithPreInit(C);
 }
 void OMPClauseEnqueue::VisitOMPDistScheduleClause(
     const OMPDistScheduleClause *C) {

Best regards,

Alexey Bataev

Software Engineer
Intel Compiler Team

[sfantao]

Hi Alexey,

I am sorry but I don't think I am following. So in your example:

#pragma omp target map(to: this->A)
[]()->void {++A}();

the map clause is going to make sure the runtime library allocates the section of this that contains A and that is what is used in the device. There are no new instances of A created inside the kernel. If you do:

#pragma omp target
[]()->void {++A}();

what the lambda is going to get would be the same (except that runtime had eventually to copy more data to the device), so the map clause has no effect in how captures are dealt with inside. All the map clause does is to inform the runtime about sizes and addresses and change the captures to byCopy when necessary, it does not create new instances of data inside the kernel. Maybe what you are proposing is meant for (first)private?

Thanks,
Samuel

[ABataev]

Samuel, this is true for device part of codegen. But what about host? If
this code must be executed on host, will it be handled correctly?

Best regards,

Alexey Bataev

Software Engineer
Intel Compiler Team

18.03.2016 23:03, Samuel Antao пишет:

sfantao added a comment.

Hi Alexey,

I am sorry but I don't think I am following. So in your example:

#pragma omp target map(to: this->A)
[]()->void {++A}();

the map clause is going to make sure the runtime library allocates the section of this that contains A and that is what is used in the device. There are no new instances of A created inside the kernel. If you do:

#pragma omp target
[]()->void {++A}();

what the lambda is going to get would be the same (except that runtime had eventually to copy more data to the device), so the map clause has no effect in how captures are dealt with inside. All the map clause does is to inform the runtime about sizes and addresses and change the captures to byCopy when necessary, it does not create new instances of data inside the kernel. Maybe what you are proposing is meant for (first)private?

Thanks,
Samuel

http://reviews.llvm.org/D16749

[sfantao]

In D16749#378969, @ABataev wrote:

Samuel, this is true for device part of codegen. But what about host? If
this code must be executed on host, will it be handled correctly?

Hi Alexey,

Yes, it will be handled correctly. The map clause doesn't change the way codegen is done for the host. The host always takes the original references regardless the map clause and its modifiers. Even for the device side, map clause won't have any effect in the code generation of the outlined target function. This behavior is what ensures that everything works seamlessly in a system that has a shared address space for host and device.

In any case, I added a new regression tests (CK25) that ensures that the lambda gets the same captures regardless the map clause.

[sfantao]

• Add regression tests for inner captures in the presence of map clause.
• Rebase.

[sfantao]

Ping!

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

3853–3891

I do recall that similar code already exists in Sema. Maybe you'd better to store info about top-level bases in map clause in Sema rather than copy the code between modules?

3922

Replace 'auto' by 'QualType'. Also, I think all such deep digging into structure of expressions must reside in Sema rather than in codegen and info must be stored within OMPMapClause

3953–3958

Do we really need this? It is quite simple and has nothing to do with the name of the function

4141–4144

CGF has EmitLoadOfPointer() function

4159–4195

Maybe it is better to extract it as a member function?

4160

auto &&

4254

getVarRefs()->varlists(). getVarRefs() is protected

[sfantao]

• Use the new mappable expression information stored in the map clause by Sema.
• Add tests for mapped fields that are privatized in the current scope.
• Rebase.

[sfantao]

Hi Alexey,

Thanks for the review!

lib/CodeGen/CGOpenMPRuntime.cpp

3853–3891

Yes, Sema does something similar. In the last diff I use what Sema produces and store it in the clause. The implementation is in the dependence diff.

3922

Done. The information is now stored in the map clause.

3953–3958

Ok, I removed this.

4141–4144

Ok, using EmitLoadOfPointerLValue now.

4159–4195

Moved this code to a new member function as suggested.

4160

I am now using a member function as you suggested below.

4254

I am doing as you say. In any case, the const variant of getVarRefs is currently public member. Maybe that is something you plan to fix?

[sfantao]

• Rebase.

[sfantao]

• Use canonical declaration when dealng with mappable expressions.

[sfantao]

• Rebase.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

3862–3864

Remove braces

4256

cast is not required, remove it

[sfantao]

• Remove extra braces.

[sfantao]

Hi Alexey,

Thanks for the review!

lib/CodeGen/CGOpenMPRuntime.cpp

3862–3864

Done!

4256

It does not work without the cast because the base class Decl returns a Decl*. Even if the overrides in the child classes use the child type, the base class is what is used for implicit cast checking.

lib/CodeGen/CGExpr.cpp

LValue CodeGenFunction::EmitOMPArraySectionExpr(const OMPArraySectionExpr *E,
                                                 bool IsLowerBound) {
   LValue Base;
   if (auto *ASE =
           dyn_cast<OMPArraySectionExpr>(E->getBase()->IgnoreParenImpCasts()))
     Base = EmitOMPArraySectionExpr(ASE, IsLowerBound);
   else
     Base = EmitLValue(E->getBase());
   QualType BaseTy = Base.getType();
+  // If the base is a pointer, we actually need to compute the offset in the
+  // pointee.
+  if (BaseTy->isAnyPointerType()) {
+    auto RBase = EmitLoadOfLValue(Base, E->getExprLoc());
+    Base = MakeNaturalAlignAddrLValue(
+        RBase.getScalarVal(), BaseTy->getAs<PointerType>()->getPointeeType());
+  }
   llvm::Value *Idx = nullptr;
   QualType ResultExprTy;
   if (auto *AT = getContext().getAsArrayType(BaseTy))
     ResultExprTy = AT->getElementType();
   else
     ResultExprTy = BaseTy->getPointeeType();
   if (IsLowerBound || (!IsLowerBound && E->getColonLoc().isInvalid())) {
     // Requesting lower bound or upper bound, but without provided length and

lib/CodeGen/CGOpenMPRuntime.cpp

OutlinedFnID = new llvm::GlobalVariable(
         llvm::Constant::getNullValue(CGM.Int8Ty), ".omp_offload.region_id");
 
   // Register the information for the entry associated with this target region.
   OffloadEntriesInfoManager.registerTargetRegionEntryInfo(
       DeviceID, FileID, ParentName, Line, Column, OutlinedFn, OutlinedFnID);
   return;
 }
 
-void CGOpenMPRuntime::emitTargetCall(CodeGenFunction &CGF,
-                                     const OMPExecutableDirective &D,
-                                     llvm::Value *OutlinedFn,
-                                     llvm::Value *OutlinedFnID,
-                                     const Expr *IfCond, const Expr *Device,
-                                     ArrayRef<llvm::Value *> CapturedVars) {
-  if (!CGF.HaveInsertPoint())
-    return;
+namespace {
+// \brief Utility to extract information from the map clauses associated with a
+// given construct and provide a convenient interface to obtain the information
+// and generate code for that information.
+class OpenMPMapClauseHandler {
+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.
     OMP_MAP_FROM = 0x02,
+    /// \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,
     /// \brief Pass the element to the device by value.
     OMP_MAP_BYCOPY = 0x80,
   };
 
+  typedef SmallVector<llvm::Value *, 16> MapValuesArrayTy;
+  typedef SmallVector<unsigned, 16> MapFlagsArrayTy;
+
+private:
+  /// \brief Directive from where the map clauses were extracted.
+  const OMPExecutableDirective &Directive;
+
+  /// \brief Function the directive is being generated for.
+  CodeGenFunction &CGF;
+
+  struct DeclarationMapInfoEntry {
+    /// \brief Array of components in the map expression.
+    typedef std::pair<const Expr *, const ValueDecl *> ComponentTy;
+    typedef llvm::SmallVector<ComponentTy, 4> ComponentsTy;
+    ComponentsTy Components;
+
+    // Map type and modifier associated with this expression.
+    OpenMPMapClauseKind MapType;
+    OpenMPMapClauseKind MapTypeModifier;
+
+    /// \brief Build and initialize this map information record with information
+    /// retrieved from the provided map clause expression.
+    DeclarationMapInfoEntry(const Expr *MCE, OpenMPMapClauseKind MapType,
+                            OpenMPMapClauseKind MapTypeModifier)
+        : MapType(MapType), MapTypeModifier(MapTypeModifier) {
+      assert(MCE && "Invalid expression??");
+      while (true) {
+        MCE = MCE->IgnoreParenImpCasts();
+
+        if (auto *CurE = dyn_cast<DeclRefExpr>(MCE)) {
+          Components.push_back(
+              ComponentTy(CurE, cast<VarDecl>(CurE->getDecl())));
+          break;
+        }
+
+        if (auto *CurE = dyn_cast<MemberExpr>(MCE)) {
+          auto *BaseE = CurE->getBase()->IgnoreParenImpCasts();
+
+          Components.push_back(
+              ComponentTy(CurE, cast<FieldDecl>(CurE->getMemberDecl())));
+          if (isa<CXXThisExpr>(BaseE))
+            break;
+
+          MCE = BaseE;
+          continue;
+        }
+
+        if (auto *CurE = dyn_cast<ArraySubscriptExpr>(MCE)) {
+          Components.push_back(ComponentTy(CurE, nullptr));
+          MCE = CurE->getBase()->IgnoreParenImpCasts();
+          continue;
+        }
+
+        if (auto *CurE = dyn_cast<OMPArraySectionExpr>(MCE)) {
+          Components.push_back(ComponentTy(CurE, nullptr));
+          MCE = CurE->getBase()->IgnoreParenImpCasts();
+          continue;
+        }
+
+        llvm_unreachable("Invalid map clause expression!");
+      }
+    }
+
+    /// \brief Return declaration associated with this map information. If it is
+    /// a field it means the base is 'this'.
+    const ValueDecl *getAssociatedDecl() const {
+      assert(!Components.empty() &&
+             "No expressions to extract declaration from??");
+      const ValueDecl *D = Components.back().second;
+      assert(D && "Declaration must exist!");
+      return D;
+    }
+  };
+
+  /// \brief Map between a declaration and its associated map information
+  /// entries. If the map info relates to 'this' we map it to null.
+  typedef SmallVector<DeclarationMapInfoEntry *, 4> DeclarationMapInfoEntriesTy;
+  llvm::DenseMap<const ValueDecl *, DeclarationMapInfoEntriesTy>
+      DeclarationMapInfoMap;
+
+  llvm::Value *getExprTypeSize(const Expr *E) const {
+    auto ExprTy = E->getType().getCanonicalType();
+
+    // Reference types are ignored for mapping purposes.
+    if (auto *RefTy = ExprTy->getAs<ReferenceType>())
+      ExprTy = RefTy->getPointeeType().getCanonicalType();
+
+    // Given that an array section is considered a built-in type, we need to
+    // do the calculation based on the length of the section instead of relying
+    // on CGF.getTypeSize(E->getType()).
+    if (const auto *OAE = dyn_cast<OMPArraySectionExpr>(E)) {
+
+      auto BaseTy = OAE->getBase()->getType().getCanonicalType();
+      // Reference types are ignored for mapping purposes.
+      if (auto *RefTy = BaseTy->getAs<ReferenceType>())
+        BaseTy = RefTy->getPointeeType().getCanonicalType();
+
+      // If there is no length associated with the expression, that means we
+      // are using the whole length of the base.
+      if (!OAE->getLength())
+        return CGF.getTypeSize(BaseTy);
+
+      llvm::Value *ElemSize;
+      if (auto *PTy = BaseTy->getAs<PointerType>()) {
+        ElemSize = CGF.getTypeSize(PTy->getPointeeType().getCanonicalType());
+      } else {
+        auto *ATy = cast<ArrayType>(BaseTy.getTypePtr());
+        assert(ATy && "Expecting array type if not a pointer type.");
+        ElemSize = CGF.getTypeSize(ATy->getElementType().getCanonicalType());
+      }
+
+      auto *LengthVal = CGF.EmitScalarExpr(OAE->getLength());
+      LengthVal =
+          CGF.Builder.CreateIntCast(LengthVal, CGF.SizeTy, /*isSigned=*/false);
+      return CGF.Builder.CreateNUWMul(LengthVal, ElemSize);
+    }
+    return CGF.getTypeSize(ExprTy);
+  }
+
+  /// \brief Generate the address of the lower bound of the section defined by
+  /// expression \a E.
+  llvm::Value *getLowerBoundOfElement(const Expr *E) const {
+    return CGF.EmitLValue(E).getPointer();
+  }
+
+  /// \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.
+  unsigned getMapTypeBits(const DeclarationMapInfoEntry *Entry, bool AddPtrFlag,
+                          bool AddExtraFlag) const {
+    unsigned Bits = 0u;
+    switch (Entry->MapType) {
+    case OMPC_MAP_alloc:
+      Bits = OMP_MAP_ALLOC;
+      break;
+    case OMPC_MAP_to:
+      Bits = OMP_MAP_TO;
+      break;
+    case OMPC_MAP_from:
+      Bits = OMP_MAP_FROM;
+      break;
+    case OMPC_MAP_tofrom:
+      Bits = OMP_MAP_TO | OMP_MAP_FROM;
+      break;
+    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;
+    if (Entry->MapTypeModifier == OMPC_MAP_always)
+      Bits |= OMP_MAP_ALWAYS;
+    return Bits;
+  }
+
+  /// \brief Generate the base pointers, section pointers, sizes and map type
+  /// bits for a given set of expressions \a MIE associated with a declaration.
+  void generateInfoForEntries(const DeclarationMapInfoEntriesTy &MIE,
+                              MapValuesArrayTy &BasePointers,
+                              MapValuesArrayTy &Pointers,
+                              MapValuesArrayTy &Sizes,
+                              MapFlagsArrayTy &Types) const {
+    // The following summarizes what has to be generated for each map and the
+    // types bellow. The generated information is expressed in this order:
+    // base pointer, section pointer, size, flags
+    // (to add to the ones that come from the map type and modifier).
+    //
+    // double d;
+    // int i[100];
+    // float *p;
+    //
+    // struct S1 {
+    //   int i;
+    //   float f[50];
+    // }
+    // struct S2 {
+    //   int i;
+    //   float f[50];
+    //   S1 s;
+    //   double *p;
+    //   struct S2 *ps;
+    // }
+    // S2 s;
+    // S2 *ps;
+    //
+    // map(d)
+    // &d, &d, sizeof(double), noflags
+    //
+    // map(i)
+    // &i, &i, 100*sizeof(int), noflags
+    //
+    // map(i[1:23])
+    // &i(=&i[0]), &i[1], 23*sizeof(int), noflags
+    //
+    // map(p)
+    // &p, &p, sizeof(float*), noflags
+    //
+    // map(p[1:24])
+    // p, &p[1], 24*sizeof(float), noflags
+    //
+    // map(s)
+    // &s, &s, sizeof(S2), noflags
+    //
+    // map(s.i)
+    // &s, &(s.i), sizeof(int), noflags
+    //
+    // map(s.s.f)
+    // &s, &(s.i.f), 50*sizeof(int), noflags
+    //
+    // map(s.p)
+    // &s, &(s.p), sizeof(double*), noflags
+    //
+    // map(s.p[:22], s.a s.b)
+    // &s, &(s.p), sizeof(double*), noflags
+    // &(s.p), &(s.p[0]), 22*sizeof(double), ptr_flag + extra_flag
+    //
+    // map(s.ps)
+    // &s, &(s.ps), sizeof(S2*), noflags
+    //
+    // map(s.ps->s.i)
+    // &s, &(s.ps), sizeof(S2*), noflags
+    // &(s.ps), &(s.ps->s.i), sizeof(int), ptr_flag + extra_flag
+    //
+    // map(s.ps->ps)
+    // &s, &(s.ps), sizeof(S2*), noflags
+    // &(s.ps), &(s.ps->ps), sizeof(S2*), ptr_flag + extra_flag
+    //
+    // map(s.ps->ps->ps)
+    // &s, &(s.ps), sizeof(S2*), noflags
+    // &(s.ps), &(s.ps->ps), sizeof(S2*), ptr_flag + extra_flag
+    // &(s.ps->ps), &(s.ps->ps->ps), sizeof(S2*), ptr_flag + extra_flag
+    //
+    // map(s.ps->ps->s.f[:22])
+    // &s, &(s.ps), sizeof(S2*), noflags
+    // &(s.ps), &(s.ps->ps), sizeof(S2*), ptr_flag + extra_flag
+    // &(s.ps->ps), &(s.ps->ps->s.f[0]), 22*sizeof(float), ptr_flag + extra_flag
+    //
+    // map(ps)
+    // &ps, &ps, sizeof(S2*), noflags
+    //
+    // map(ps->i)
+    // ps, &(ps->i), sizeof(int), noflags
+    //
+    // map(ps->s.f)
+    // ps, &(ps->s.f[0]), 50*sizeof(float), noflags
+    //
+    // map(ps->p)
+    // ps, &(ps->p), sizeof(double*), noflags
+    //
+    // map(ps->p[:22])
+    // ps, &(ps->p), sizeof(double*), noflags
+    // &(ps->p), &(ps->p[0]), 22*sizeof(double), ptr_flag + extra_flag
+    //
+    // map(ps->ps)
+    // ps, &(ps->ps), sizeof(S2*), noflags
+    //
+    // map(ps->ps->s.i)
+    // ps, &(ps->ps), sizeof(S2*), noflags
+    // &(ps->ps), &(ps->ps->s.i), sizeof(int), ptr_flag + extra_flag
+    //
+    // map(ps->ps->ps)
+    // ps, &(ps->ps), sizeof(S2*), noflags
+    // &(ps->ps), &(ps->ps->ps), sizeof(S2*), ptr_flag + extra_flag
+    //
+    // map(ps->ps->ps->ps)
+    // ps, &(ps->ps), sizeof(S2*), noflags
+    // &(ps->ps), &(ps->ps->ps), sizeof(S2*), ptr_flag + extra_flag
+    // &(ps->ps->ps), &(ps->ps->ps->ps), sizeof(S2*), ptr_flag + extra_flag
+    //
+    // map(ps->ps->ps->s.f[:22])
+    // ps, &(ps->ps), sizeof(S2*), noflags
+    // &(ps->ps), &(ps->ps->ps), sizeof(S2*), ptr_flag + extra_flag
+    // &(ps->ps->ps), &(ps->ps->ps->s.f[0]), 22*sizeof(float), ptr_flag +
+    // extra_flag
+
+    bool IsEntriesFirstInfo = true;
+
+    // For each expression in the map clauses...
+    for (auto *InfoForExpr : MIE) {
+      auto CI = InfoForExpr->Components.rbegin();
+      auto CE = InfoForExpr->Components.rend();
+      auto I = CI;
+
+      bool IsExpressionFirstInfo = true;
+      llvm::Value *BP = nullptr;
+
+      if (auto *ME = dyn_cast<MemberExpr>(I->first)) {
+        // The base is the 'this' pointer. The content of the pointer is going
+        // to be the base of the field being mapped.
+        BP = CGF.EmitScalarExpr(ME->getBase());
+      } else {
+        // The base is the reference to the variable.
+        // BP = &Var.
+        BP = CGF.EmitLValue(cast<DeclRefExpr>(I->first)).getPointer();
+
+        // If the variable is a pointer and is being dereferenced (i.e. is not
+        // the last component), the base has to be the pointer itself, not his
+        // reference.
+        if (I->second->getType()->isAnyPointerType() && std::next(I) != CE) {
+          auto PtrAddr =
+              CGF.MakeNaturalAlignAddrLValue(BP, I->second->getType());
+          BP = CGF.EmitLoadOfLValue(PtrAddr, SourceLocation()).getScalarVal();
+
+          // We do not need to generate individual map information for the
+          // pointer, it can be associated with the combined storage.
+          ++I;
+        }
+      }
+
+      for (; I != CE; ++I) {
+        auto Next = std::next(I);
+
+        // We need to generate the addresses and sizes if this is the last
+        // component, or if the component is a pointer. In this case, the
+        // pointer becomes the base address for the following components.
+        if (Next == CE || I->first->getType()->isAnyPointerType()) {
+
+          // If this is not the last component, we expect the pointer to be
+          // associated with an array expression or member expression.
+          assert((Next == CE || isa<MemberExpr>(Next->first) ||
+                  isa<ArraySubscriptExpr>(Next->first) ||
+                  isa<OMPArraySectionExpr>(Next->first)) &&
+                 "Unexpected expression");
+
+          // Save the base we are currently using.
+          BasePointers.push_back(BP);
+
+          auto *LB = getLowerBoundOfElement(I->first);
+          auto *Size = getExprTypeSize(I->first);
+
+          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).
+          Types.push_back(getMapTypeBits(InfoForExpr, !IsExpressionFirstInfo,
+                                         !IsEntriesFirstInfo));
+
+          // The pointer becomes the base for the next element.
+          if (Next != CE)
+            BP = LB;
+
+          IsExpressionFirstInfo = false;
+          IsEntriesFirstInfo = false;
+          continue;
+        }
+      }
+    }
+    return;
+  }
+
+public:
+  OpenMPMapClauseHandler(const OMPExecutableDirective &Dir,
+                         CodeGenFunction &CGF)
+      : Directive(Dir), CGF(CGF) {
+
+    // Scan and extract information from the map clauses one by one.
+    for (auto *MC : Directive.getClausesOfKind<OMPMapClause>()) {
+      for (auto *RE : MC->getVarRefs()) {
+        auto *MI = new DeclarationMapInfoEntry(RE, MC->getMapType(),
+                                               MC->getMapTypeModifier());
+        const auto *D = MI->getAssociatedDecl();
+        if (isa<FieldDecl>(D))
+          DeclarationMapInfoMap[nullptr].push_back(MI);
+        else
+          DeclarationMapInfoMap[D].push_back(MI);
+      }
+    }
+  }
+
+  ~OpenMPMapClauseHandler() {
+    // Clear the entries.
+    for (auto &MapEntries : DeclarationMapInfoMap)
+      for (auto *MapEntry : MapEntries.second)
+        delete (MapEntry);
+  }
+
+  /// \brief Generate all the base pointers, section pointers, sizes and map
+  /// types for the extracted map information.
+  void generateAllInfo(MapValuesArrayTy &BasePointers,
+                       MapValuesArrayTy &Pointers, MapValuesArrayTy &Sizes,
+                       MapFlagsArrayTy &Types) const {
+    BasePointers.clear();
+    Pointers.clear();
+    Sizes.clear();
+    Types.clear();
+
+    // For each declaration identified in the map clause...
+    for (auto &MapEntries : DeclarationMapInfoMap) {
+      const DeclarationMapInfoEntriesTy &MIE = MapEntries.second;
+      generateInfoForEntries(MIE, BasePointers, Pointers, Sizes, Types);
+    }
+    return;
+  }
+
+  /// \brief Generate the base pointers, section pointers, sizes and map types
+  /// associated to a given capture.
+  void generateInfoForCapture(const CapturedStmt::Capture *Cap,
+                              MapValuesArrayTy &BasePointers,
+                              MapValuesArrayTy &Pointers,
+                              MapValuesArrayTy &Sizes,
+                              MapFlagsArrayTy &Types) const {
+    assert(!Cap->capturesVariableArrayType() &&
+           "Not expecting to generate map info for a variable array type!");
+
+    BasePointers.clear();
+    Pointers.clear();
+    Sizes.clear();
+    Types.clear();
+
+    const ValueDecl *VD = Cap->capturesThis() ? nullptr : Cap->getCapturedVar();
+    auto I = DeclarationMapInfoMap.find(VD);
+    if (I == DeclarationMapInfoMap.end())
+      return;
+
+    const DeclarationMapInfoEntriesTy &MIE = I->second;
+    assert(!MIE.empty() && "Not expecting declaration with empty information!");
+    generateInfoForEntries(MIE, BasePointers, Pointers, Sizes, Types);
+
+    return;
+  }
+};
+}
+
+void CGOpenMPRuntime::emitTargetCall(CodeGenFunction &CGF,
+                                     const OMPExecutableDirective &D,
+                                     llvm::Value *OutlinedFn,
+                                     llvm::Value *OutlinedFnID,
+                                     const Expr *IfCond, const Expr *Device,
+                                     ArrayRef<llvm::Value *> CapturedVars) {
+  if (!CGF.HaveInsertPoint())
+    return;
+
   enum OpenMPOffloadingReservedDeviceIDs {
     /// \brief Device ID if the device was not defined, runtime should get it
     /// from environment variables in the spec.
     OMP_DEVICEID_UNDEF = -1,
   };
 
   assert(OutlinedFn && "Invalid outlined function!");
 
   auto &Ctx = CGF.getContext();
 
-  // Fill up the arrays with the all the captured variables.
-  SmallVector<llvm::Value *, 16> BasePointers;
-  SmallVector<llvm::Value *, 16> Pointers;
-  SmallVector<llvm::Value *, 16> Sizes;
-  SmallVector<unsigned, 16> MapTypes;
+  // Fill up the arrays with all the captured variables.
+  OpenMPMapClauseHandler::MapValuesArrayTy KernelArgs;
+  OpenMPMapClauseHandler::MapValuesArrayTy BasePointers;
+  OpenMPMapClauseHandler::MapValuesArrayTy Pointers;
+  OpenMPMapClauseHandler::MapValuesArrayTy Sizes;
+  OpenMPMapClauseHandler::MapFlagsArrayTy MapTypes;
+
+  OpenMPMapClauseHandler::MapValuesArrayTy CurBasePointers;
+  OpenMPMapClauseHandler::MapValuesArrayTy CurPointers;
+  OpenMPMapClauseHandler::MapValuesArrayTy CurSizes;
+  OpenMPMapClauseHandler::MapFlagsArrayTy CurMapTypes;
 
-  bool hasVLACaptures = false;
+  // Get map clause information.
+  OpenMPMapClauseHandler MCHandler(D, CGF);
 
   const CapturedStmt &CS = *cast<CapturedStmt>(D.getAssociatedStmt());
   auto RI = CS.getCapturedRecordDecl()->field_begin();
-  // auto II = CS.capture_init_begin();
   auto CV = CapturedVars.begin();
   for (CapturedStmt::const_capture_iterator CI = CS.capture_begin(),
                                             CE = CS.capture_end();
        CI != CE; ++CI, ++RI, ++CV) {
     StringRef Name;
     QualType Ty;
-    llvm::Value *BasePointer;
-    llvm::Value *Pointer;
-    llvm::Value *Size;
-    unsigned MapType;
 
-    // VLA sizes are passed to the outlined region by copy.
+    CurBasePointers.clear();
+    CurPointers.clear();
+    CurSizes.clear();
+    CurMapTypes.clear();
+
+    // VLA sizes are passed to the outlined region by copy and do not have map
+    // information associated.
     if (CI->capturesVariableArrayType()) {
-      BasePointer = Pointer = *CV;
-      Size = CGF.getTypeSize(RI->getType());
+      CurBasePointers.push_back(*CV);
+      CurPointers.push_back(*CV);
+      CurSizes.push_back(CGF.getTypeSize(RI->getType()));
       // Copy to the device as an argument. No need to retrieve it.
-      MapType = OMP_MAP_BYCOPY;
-      hasVLACaptures = true;
-    } else if (CI->capturesThis()) {
-      BasePointer = Pointer = *CV;
-      const PointerType *PtrTy = cast<PointerType>(RI->getType().getTypePtr());
-      Size = CGF.getTypeSize(PtrTy->getPointeeType());
+      CurMapTypes.push_back(OpenMPMapClauseHandler::OMP_MAP_BYCOPY);
+    } else {
+      // If we have any information in the map clause, we use it, otherwise we
+      // just do a default mapping.
+      MCHandler.generateInfoForCapture(CI, CurBasePointers, CurPointers,
+                                       CurSizes, CurMapTypes);
+
+      if (CurBasePointers.empty()) {
+        // Do the default mapping.
+        if (CI->capturesThis()) {
+          CurBasePointers.push_back(*CV);
+          CurPointers.push_back(*CV);
+          const PointerType *PtrTy =
+              cast<PointerType>(RI->getType().getTypePtr());
+          CurSizes.push_back(CGF.getTypeSize(PtrTy->getPointeeType()));
           // Default map type.
-      MapType = OMP_MAP_TO | OMP_MAP_FROM;
+          CurMapTypes.push_back(OpenMPMapClauseHandler::OMP_MAP_TO |
+                                OpenMPMapClauseHandler::OMP_MAP_FROM);
         } else if (CI->capturesVariableByCopy()) {
-      MapType = OMP_MAP_BYCOPY;
+          CurMapTypes.push_back(OpenMPMapClauseHandler::OMP_MAP_BYCOPY);
           if (!RI->getType()->isAnyPointerType()) {
-        // If the field is not a pointer, we need to save the actual value and
+            // If the field is not a pointer, we need to save the actual value
+            // and
             // load it as a void pointer.
             auto DstAddr = CGF.CreateMemTemp(
                 Ctx.getUIntPtrType(),
                 Twine(CI->getCapturedVar()->getName()) + ".casted");
             LValue DstLV = CGF.MakeAddrLValue(DstAddr, Ctx.getUIntPtrType());
 
             auto *SrcAddrVal = CGF.EmitScalarConversion(
                 DstAddr.getPointer(), Ctx.getPointerType(Ctx.getUIntPtrType()),
                 Ctx.getPointerType(RI->getType()), SourceLocation());
             LValue SrcLV =
                 CGF.MakeNaturalAlignAddrLValue(SrcAddrVal, RI->getType());
 
             // Store the value using the source type pointer.
             CGF.EmitStoreThroughLValue(RValue::get(*CV), SrcLV);
 
             // Load the value using the destination type pointer.
-        BasePointer = Pointer =
-            CGF.EmitLoadOfLValue(DstLV, SourceLocation()).getScalarVal();
+            CurBasePointers.push_back(
+                CGF.EmitLoadOfLValue(DstLV, SourceLocation()).getScalarVal());
+            CurPointers.push_back(CurBasePointers.back());
           } else {
-        MapType |= OMP_MAP_PTR;
-        BasePointer = Pointer = *CV;
+            CurBasePointers.push_back(*CV);
+            CurPointers.push_back(*CV);
           }
-      Size = CGF.getTypeSize(RI->getType());
+          CurSizes.push_back(CGF.getTypeSize(RI->getType()));
         } else {
           assert(CI->capturesVariable() && "Expected captured reference.");
-      BasePointer = Pointer = *CV;
+          CurBasePointers.push_back(*CV);
+          CurPointers.push_back(*CV);
 
           const ReferenceType *PtrTy =
               cast<ReferenceType>(RI->getType().getTypePtr());
           QualType ElementType = PtrTy->getPointeeType();
-      Size = CGF.getTypeSize(ElementType);
+          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,
+          // default the value doesn't have to be retrieved. For an aggregate
+          // type,
           // the default is 'tofrom'.
-      MapType = ElementType->isAggregateType() ? (OMP_MAP_TO | OMP_MAP_FROM)
-                                               : OMP_MAP_TO;
-      if (ElementType->isAnyPointerType())
-        MapType |= OMP_MAP_PTR;
+          CurMapTypes.push_back(ElementType->isAggregateType()
+                                    ? (OpenMPMapClauseHandler::OMP_MAP_TO |
+                                       OpenMPMapClauseHandler::OMP_MAP_FROM)
+                                    : OpenMPMapClauseHandler::OMP_MAP_TO);
         }
-
-    BasePointers.push_back(BasePointer);
-    Pointers.push_back(Pointer);
-    Sizes.push_back(Size);
-    MapTypes.push_back(MapType);
+      }
+    }
+    // We expect to have at least an element of information for this capture.
+    assert(!CurBasePointers.empty() && "Non-existing map pointer for capture!");
+    assert(CurBasePointers.size() == CurPointers.size() &&
+           CurBasePointers.size() == CurSizes.size() &&
+           CurBasePointers.size() == CurMapTypes.size() &&
+           "Inconsistent map information sizes!");
+
+    // The kernel args are always the first elements of the base pointers
+    // associated with a capture.
+    KernelArgs.push_back(CurBasePointers.front());
+    // We need to append the results of this capture to what we already have.
+    BasePointers.append(CurBasePointers.begin(), CurBasePointers.end());
+    Pointers.append(CurPointers.begin(), CurPointers.end());
+    Sizes.append(CurSizes.begin(), CurSizes.end());
+    MapTypes.append(CurMapTypes.begin(), CurMapTypes.end());
+  }
+
+  // Detect if we have any capture size requiring runtime evaluation of the size
+  // so that a constant array could be eventually used.
+  bool hasRuntimeEvaluationCaptureSize = false;
+  for (auto *S : Sizes)
+    if (!isa<llvm::Constant>(S)) {
+      hasRuntimeEvaluationCaptureSize = true;
+      break;
     }
 
   // Keep track on whether the host function has to be executed.
   auto OffloadErrorQType =
       Ctx.getIntTypeForBitwidth(/*DestWidth=*/32, /*Signed=*/true);
   auto OffloadError = CGF.MakeAddrLValue(
       CGF.CreateMemTemp(OffloadErrorQType, ".run_host_version"),
       OffloadErrorQType);
   CGF.EmitStoreOfScalar(llvm::Constant::getNullValue(CGM.Int32Ty),
                         OffloadError);
 
   // Fill up the pointer arrays and transfer execution to the device.
   auto &&ThenGen = [this, &Ctx, &BasePointers, &Pointers, &Sizes, &MapTypes,
-                    hasVLACaptures, Device, OutlinedFnID, OffloadError,
-                    OffloadErrorQType](CodeGenFunction &CGF) {
+                    hasRuntimeEvaluationCaptureSize, Device, OutlinedFnID,
+                    OffloadError, OffloadErrorQType](CodeGenFunction &CGF) {
     unsigned PointerNumVal = BasePointers.size();
     llvm::Value *PointerNum = CGF.Builder.getInt32(PointerNumVal);
     llvm::Value *BasePointersArray;
     llvm::Value *PointersArray;
     llvm::Value *SizesArray;
     llvm::Value *MapTypesArray;
 
     if (PointerNumVal) {
       llvm::APInt PointerNumAP(32, PointerNumVal, /*isSigned=*/true);
       QualType PointerArrayType = Ctx.getConstantArrayType(
           Ctx.VoidPtrTy, PointerNumAP, ArrayType::Normal,
           /*IndexTypeQuals=*/0);
 
       BasePointersArray =
           CGF.CreateMemTemp(PointerArrayType, ".offload_baseptrs").getPointer();
       PointersArray =
           CGF.CreateMemTemp(PointerArrayType, ".offload_ptrs").getPointer();
 
       // If we don't have any VLA types, we can use a constant array for the map
       // sizes, otherwise we need to fill up the arrays as we do for the
       // pointers.
-      if (hasVLACaptures) {
+      if (hasRuntimeEvaluationCaptureSize) {
         QualType SizeArrayType = Ctx.getConstantArrayType(
             Ctx.getSizeType(), PointerNumAP, ArrayType::Normal,
             /*IndexTypeQuals=*/0);
         SizesArray =
             CGF.CreateMemTemp(SizeArrayType, ".offload_sizes").getPointer();
       } else {
         // We expect all the sizes to be constant, so we collect them to create
         // a constant array.
if (PointerNumVal) {
           PVal = CGF.Builder.CreateIntToPtr(PVal, CGM.VoidPtrTy);
         }
         llvm::Value *P = CGF.Builder.CreateConstInBoundsGEP2_32(
             llvm::ArrayType::get(CGM.VoidPtrTy, PointerNumVal), PointersArray,
             0, i);
         Address PAddr(P, Ctx.getTypeAlignInChars(Ctx.VoidPtrTy));
         CGF.Builder.CreateStore(PVal, PAddr);
 
-        if (hasVLACaptures) {
+        if (hasRuntimeEvaluationCaptureSize) {
           llvm::Value *S = CGF.Builder.CreateConstInBoundsGEP2_32(
               llvm::ArrayType::get(CGM.SizeTy, PointerNumVal), SizesArray,
               /*Idx0=*/0,
               /*Idx1=*/i);
           Address SAddr(S, Ctx.getTypeAlignInChars(Ctx.getSizeType()));
           CGF.Builder.CreateStore(CGF.Builder.CreateIntCast(
                                       Sizes[i], CGM.SizeTy, /*isSigned=*/true),
                                   SAddr);
void CGOpenMPRuntime::emitTargetCall(CodeGenFunction &CGF,
   // Check the error code and execute the host version if required.
   auto OffloadFailedBlock = CGF.createBasicBlock("omp_offload.failed");
   auto OffloadContBlock = CGF.createBasicBlock("omp_offload.cont");
   auto OffloadErrorVal = CGF.EmitLoadOfScalar(OffloadError, SourceLocation());
   auto Failed = CGF.Builder.CreateIsNotNull(OffloadErrorVal);
   CGF.Builder.CreateCondBr(Failed, OffloadFailedBlock, OffloadContBlock);
 
   CGF.EmitBlock(OffloadFailedBlock);
-  CGF.Builder.CreateCall(OutlinedFn, BasePointers);
+  CGF.Builder.CreateCall(OutlinedFn, KernelArgs);
   CGF.EmitBranch(OffloadContBlock);
 
   CGF.EmitBlock(OffloadContBlock, /*IsFinished=*/true);
   return;
 }
 
 void CGOpenMPRuntime::scanForTargetRegionsFunctions(const Stmt *S,
                                                     StringRef ParentName) {

lib/Sema/SemaOpenMP.cpp

 }
 
 void Sema::InitDataSharingAttributesStack() {
   VarDataSharingAttributesStack = new DSAStackTy(*this);
 }
 
 #define DSAStack static_cast<DSAStackTy *>(VarDataSharingAttributesStack)
 
+namespace {
+// Types used to organize the components of a valid map clause.
+typedef std::pair<Expr *, ValueDecl *> MapExpressionComponent;
+typedef SmallVector<MapExpressionComponent, 4> MapExpressionComponents;
+}
+
+// Helper to extract the components in the map clause expression \a E and store
+// them into \a MEC. This assumes that \a E is a valid map clause expression,
+// i.e. it has already passed the single clause checks.
+static void ExtractMapExpressionComponents(Expr *TE,
+                                           MapExpressionComponents &MEC) {
+  while (true) {
+    TE = TE->IgnoreParenImpCasts();
+
+    if (auto *CurE = dyn_cast<DeclRefExpr>(TE)) {
+      MEC.push_back(
+          MapExpressionComponent(CurE, cast<VarDecl>(CurE->getDecl())));
+      break;
+    }
+
+    if (auto *CurE = dyn_cast<MemberExpr>(TE)) {
+      auto *BaseE = CurE->getBase()->IgnoreParenImpCasts();
+
+      MEC.push_back(
+          MapExpressionComponent(CurE, cast<FieldDecl>(CurE->getMemberDecl())));
+      if (isa<CXXThisExpr>(BaseE))
+        break;
+
+      TE = BaseE;
+      continue;
+    }
+
+    if (auto *CurE = dyn_cast<ArraySubscriptExpr>(TE)) {
+      MEC.push_back(MapExpressionComponent(CurE, nullptr));
+      TE = CurE->getBase()->IgnoreParenImpCasts();
+      continue;
+    }
+
+    if (auto *CurE = dyn_cast<OMPArraySectionExpr>(TE)) {
+      MEC.push_back(MapExpressionComponent(CurE, nullptr));
+      TE = CurE->getBase()->IgnoreParenImpCasts();
+      continue;
+    }
+
+    llvm_unreachable(
+        "Expecting only valid map clause expressions at this point!");
+  }
+}
+
 bool Sema::IsOpenMPCapturedByRef(ValueDecl *D,
                                  const CapturedRegionScopeInfo *RSI) {
   assert(LangOpts.OpenMP && "OpenMP is not allowed");
 
   auto &Ctx = getASTContext();
   bool IsByRef = true;
 
   // Find the directive that is associated with the provided scope.
if (isOpenMPTargetDirective(DKind)) {
     //    by reference, because they may have been mapped in the enclosing data
     //    environment.
     //  - If the scalar value does not fit the size of uintptr, it has to be
     //    passed by reference, regardless the result in the table above.
     //  - For pointers mapped by value that have either an implicit map or an
     //    array section, the runtime library may pass the NULL value to the
     //    device instead of the value passed to it by the compiler.
 
-    // FIXME: Right now, only implicit maps are implemented. Properly mapping
-    // values requires having the map, private, and firstprivate clauses SEMA
-    // and parsing in place, which we don't yet.
 
     if (Ty->isReferenceType())
       Ty = Ty->castAs<ReferenceType>()->getPointeeType();
+
+    // FIXME: Right now, only some of the maps are implemented. Properly mapping
+    // values requires having the private, and firstprivate clauses SEMA
+    // and parsing in place, which we don't yet.
+
+    // Locate map clauses and see if the variable being captured is referred to
+    // in any of those clauses. Here we only care about variables, not fields,
+    // because fields are part of aggregates.
+    bool IsVariableUsedInMapClause = false;
+    bool IsVariableAssociatedWithSection = false;
+
+    DSAStack->checkMapInfoForVar(
+        D, /*CurrentRegionOnly=*/true, [&](Expr *MapExpr) {
+          MapExpressionComponents MapExprComponents;
+          ExtractMapExpressionComponents(MapExpr, MapExprComponents);
+
+          auto EI = MapExprComponents.rbegin();
+          auto EE = MapExprComponents.rend();
+
+          assert(EI != EE && "Invalid map expression!");
+
+          if (isa<DeclRefExpr>(EI->first))
+            IsVariableUsedInMapClause |= EI->second == D;
+
+          ++EI;
+          if (EI == EE)
+            return false;
+
+          if (isa<ArraySubscriptExpr>(EI->first) ||
+              isa<OMPArraySectionExpr>(EI->first) ||
+              isa<MemberExpr>(EI->first)) {
+            IsVariableAssociatedWithSection = true;
+            // There is nothing more we need to know about this variable.
+            return true;
+          }
+
+          // Keep looking for more map info.
+          return false;
+        });
+
+    if (IsVariableUsedInMapClause) {
+      // If variable is identified in a map clause it is always captured by
+      // reference except if it is a pointer that is dereferenced somehow.
+      IsByRef = !(Ty->isPointerType() && IsVariableAssociatedWithSection);
+    } else {
+      // By default, all the data that has a scalar type is mapped by copy.
       IsByRef = !Ty->isScalarType();
     }
+  }
 
-  // When passing data by value, we need to make sure it fits the uintptr size
+  // When passing data by copy, we need to make sure it fits the uintptr size
   // and alignment, because the runtime library only deals with uintptr types.
   // If it does not fit the uintptr size, we need to pass the data by reference
   // instead.
   if (!IsByRef &&
       (Ctx.getTypeSizeInChars(Ty) >
            Ctx.getTypeSizeInChars(Ctx.getUIntPtrType()) ||
        Ctx.getDeclAlign(D) > Ctx.getTypeAlignInChars(Ctx.getUIntPtrType())))
     IsByRef = true;
static Expr *CheckMapClauseExpressionBase(Sema &SemaRef, Expr *E) {
   return RelevantExpr;
 }
 
 // Return true if expression E associated with value VD has conflicts with other
 // map information.
 static bool CheckMapConflicts(Sema &SemaRef, DSAStackTy *DSAS, ValueDecl *VD,
                               Expr *E, bool CurrentRegionOnly) {
   assert(VD && E);
-
-  // Types used to organize the components of a valid map clause.
-  typedef std::pair<Expr *, ValueDecl *> MapExpressionComponent;
-  typedef SmallVector<MapExpressionComponent, 4> MapExpressionComponents;
-
-  // Helper to extract the components in the map clause expression E and store
-  // them into MEC. This assumes that E is a valid map clause expression, i.e.
-  // it has already passed the single clause checks.
-  auto ExtractMapExpressionComponents = [](Expr *TE,
-                                           MapExpressionComponents &MEC) {
-    while (true) {
-      TE = TE->IgnoreParenImpCasts();
-
-      if (auto *CurE = dyn_cast<DeclRefExpr>(TE)) {
-        MEC.push_back(
-            MapExpressionComponent(CurE, cast<VarDecl>(CurE->getDecl())));
-        break;
-      }
-
-      if (auto *CurE = dyn_cast<MemberExpr>(TE)) {
-        auto *BaseE = CurE->getBase()->IgnoreParenImpCasts();
-
-        MEC.push_back(MapExpressionComponent(
-            CurE, cast<FieldDecl>(CurE->getMemberDecl())));
-        if (isa<CXXThisExpr>(BaseE))
-          break;
-
-        TE = BaseE;
-        continue;
-      }
-
-      if (auto *CurE = dyn_cast<ArraySubscriptExpr>(TE)) {
-        MEC.push_back(MapExpressionComponent(CurE, nullptr));
-        TE = CurE->getBase()->IgnoreParenImpCasts();
-        continue;
-      }
-
-      if (auto *CurE = dyn_cast<OMPArraySectionExpr>(TE)) {
-        MEC.push_back(MapExpressionComponent(CurE, nullptr));
-        TE = CurE->getBase()->IgnoreParenImpCasts();
-        continue;
-      }
-
-      llvm_unreachable(
-          "Expecting only valid map clause expressions at this point!");
-    }
-  };
-
   SourceLocation ELoc = E->getExprLoc();
   SourceRange ERange = E->getSourceRange();
 
   // In order to easily check the conflicts we need to match each component of
   // the expression under test with the components of the expressions that are
   // already in the stack.
 
   MapExpressionComponents CurComponents;

test/OpenMP/target_map_codegen.cpp

 // RUN: %clang_cc1 -DCK10 -fopenmp -omptargets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -omptargets=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 CK10
 // RUN: %clang_cc1 -DCK10 -verify -fopenmp -omptargets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK10
 // RUN: %clang_cc1 -DCK10 -fopenmp -omptargets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -omptargets=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 | OMP_MAP_PTR = 128 + 32
-// CK10-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 160]
+// Map types: OMP_MAP_BYCOPY = 128
+// CK10-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128]
 
 // CK10-LABEL: implicit_maps_pointer
 void implicit_maps_pointer (){
   double *ddyn;
 
   // CK10-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
   // CK10-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
   // CK10-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0

 // CK18: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]])
 // CK18: [[ADDR:%.+]] = alloca i[[sz]],
 // CK18: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR]],
 // CK18-64: [[CADDR:%.+]] = bitcast i64* [[ADDR]] to i32*
 // CK18-64: {{.+}} = load i32, i32* [[CADDR]],
 // CK18-32: {{.+}} = load i32, i32* [[ADDR]],
 
 #endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK19 -verify -fopenmp -omptargets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK19 --check-prefix CK19-64
+// RUN: %clang_cc1 -DCK19 -fopenmp -omptargets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -omptargets=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 CK19 --check-prefix CK19-64
+// RUN: %clang_cc1 -DCK19 -verify -fopenmp -omptargets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK19 --check-prefix CK19-32
+// RUN: %clang_cc1 -DCK19 -fopenmp -omptargets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -omptargets=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 CK19 --check-prefix CK19-32
+#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: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400]
+// CK19: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 1]
+
+// CK19: [[SIZE02:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 240]
+// CK19: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 2]
+
+// CK19: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 240]
+// CK19: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// CK19: [[SIZE04:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400]
+// CK19: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] zeroinitializer
+
+// CK19: [[SIZE05:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
+// CK19: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i32] [i32 1]
+
+// CK19: [[MTYPE06:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// CK19: [[MTYPE07:@.+]] = private {{.*}}constant [1 x i32] zeroinitializer
+
+// CK19: [[SIZE08:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
+// CK19: [[MTYPE08:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// CK19: [[SIZE09:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{8|4}}]
+// CK19: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i32] [i32 2]
+
+// CK19: [[SIZE10:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 240]
+// CK19: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// CK19: [[SIZE11:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 240]
+// CK19: [[MTYPE11:@.+]] = private {{.*}}constant [1 x i32] zeroinitializer
+
+// CK19: [[SIZE12:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
+// CK19: [[MTYPE12:@.+]] = private {{.*}}constant [1 x i32] [i32 1]
+
+// CK19: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i32] zeroinitializer
+
+// CK19: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i32] [i32 1]
+
+// CK19: [[SIZE15:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
+// CK19: [[MTYPE15:@.+]] = private {{.*}}constant [1 x i32] [i32 2]
+
+// CK19: [[MTYPE16:@.+]] = private {{.*}}constant [2 x i32] [i32 128, i32 1]
+
+// 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: [[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: [[MTYPE19:@.+]] = private {{.*}}constant [2 x i32] [i32 128, i32 0]
+
+// 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: [[MTYPE21:@.+]] = private {{.*}}constant [2 x i32] [i32 128, i32 3]
+
+// 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: [[SIZE23:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
+// CK19: [[MTYPE23:@.+]] = private {{.*}}constant [1 x i32] [i32 7]
+
+// CK19-LABEL: explicit_maps_single
+void explicit_maps_single (int ii){
+
+  // Map of a scalar.
+  int a = ii;
+
+  // Region 00
+  // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+  // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK19-DAG: [[CBPVAL0]] = bitcast i32* [[VAR0:%.+]] to i8*
+  // CK19-DAG: [[CPVAL0]] = bitcast i32* [[VAR0]] to i8*
+
+  // CK19: call void [[CALL00:@.+]](i32* {{[^,]+}})
+  #pragma omp target map(alloc:a)
+  {
+    ++a;
+  }
+
+  // Map of an array.
+  int arra[100];
+
+  // Region 01
+  // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
+  // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK19-DAG: [[CBPVAL0]] = bitcast [100 x i32]* [[VAR0:%.+]] to i8*
+  // CK19-DAG: [[CPVAL0]] = bitcast [100 x i32]* [[VAR0]] to i8*
+
+  // CK19: call void [[CALL01:@.+]]([100 x i32]* {{[^,]+}})
+  #pragma omp target map(to:arra)
+  {
+    arra[50]++;
+  }
+
+  // Region 02
+  // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}})
+  // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK19-DAG: [[CBPVAL0]] = bitcast [100 x i32]* [[VAR0:%.+]] to i8*
+  // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
+  // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[100 x i32]* [[VAR0]], i{{.+}} 0, i{{.+}} 20
+
+  // CK19: call void [[CALL02:@.+]]([100 x i32]* {{[^,]+}})
+  #pragma omp target map(from:arra[20:60])
+  {
+    arra[50]++;
+  }
+
+  // Region 03
+  // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE03]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}})
+  // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK19-DAG: [[CBPVAL0]] = bitcast [100 x i32]* [[VAR0:%.+]] to i8*
+  // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
+  // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[100 x i32]* [[VAR0]], i{{.+}} 0, i{{.+}} 0
+
+  // CK19: call void [[CALL03:@.+]]([100 x i32]* {{[^,]+}})
+  #pragma omp target map(tofrom:arra[:60])
+  {
+    arra[50]++;
+  }
+
+  // Region 04
+  // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE04]]{{.+}})
+  // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK19-DAG: [[CBPVAL0]] = bitcast [100 x i32]* [[VAR0:%.+]] to i8*
+  // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
+  // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[100 x i32]* [[VAR0]], i{{.+}} 0, i{{.+}} 0
+
+  // CK19: call void [[CALL04:@.+]]([100 x i32]* {{[^,]+}})
+  #pragma omp target map(alloc:arra[:])
+  {
+    arra[50]++;
+  }
+
+  // Region 05
+  // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE05]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE05]]{{.+}})
+  // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK19-DAG: [[CBPVAL0]] = bitcast [100 x i32]* [[VAR0:%.+]] to i8*
+  // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
+  // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[100 x i32]* [[VAR0]], i{{.+}} 0, i{{.+}} 15
+
+  // CK19: call void [[CALL05:@.+]]([100 x i32]* {{[^,]+}})
+  #pragma omp target map(to:arra[15])
+  {
+    arra[15]++;
+  }
+
+  // Region 06
+  // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE06]]{{.+}})
+  // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+  // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+
+  // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK19-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], i{{.+}}* [[S0]]
+  // CK19-DAG: [[CBPVAL0]] = bitcast [100 x i32]* [[VAR0:%.+]] to i8*
+  // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
+  // CK19-DAG: [[CSVAL0]] = mul nuw i{{.+}} %{{.*}}, 4
+  // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[100 x i32]* [[VAR0]], i{{.+}} 0, i{{.+}} %{{.*}}
+
+  // CK19: call void [[CALL06:@.+]]([100 x i32]* {{[^,]+}})
+  #pragma omp target map(tofrom:arra[ii:ii+23])
+  {
+    arra[50]++;
+  }
+
+  // Region 07
+  // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE07]]{{.+}})
+  // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+  // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+
+  // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK19-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], i{{.+}}* [[S0]]
+  // CK19-DAG: [[CBPVAL0]] = bitcast [100 x i32]* [[VAR0:%.+]] to i8*
+  // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
+  // CK19-DAG: [[CSVAL0]] = mul nuw i{{.+}} %{{.*}}, 4
+  // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[100 x i32]* [[VAR0]], i{{.+}} 0, i{{.+}} 0
+
+  // CK19: call void [[CALL07:@.+]]([100 x i32]* {{[^,]+}})
+  #pragma omp target map(alloc:arra[:ii])
+  {
+    arra[50]++;
+  }
+
+  // Region 08
+  // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE08]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE08]]{{.+}})
+  // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK19-DAG: [[CBPVAL0]] = bitcast [100 x i32]* [[VAR0:%.+]] to i8*
+  // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
+  // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[100 x i32]* [[VAR0]], i{{.+}} 0, i{{.+}} %{{.*}}
+
+  // CK19: call void [[CALL08:@.+]]([100 x i32]* {{[^,]+}})
+  #pragma omp target map(tofrom:arra[ii])
+  {
+    arra[15]++;
+  }
+
+  // Map of a pointer.
+  int *pa;
+
+  // Region 09
+  // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE09]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE09]]{{.+}})
+  // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK19-DAG: [[CBPVAL0]] = bitcast i32** [[VAR0:%.+]] to i8*
+  // CK19-DAG: [[CPVAL0]] = bitcast i32** [[VAR0]] to i8*
+
+  // CK19: call void [[CALL09:@.+]](i32** {{[^,]+}})
+  #pragma omp target map(from:pa)
+  {
+    pa[50]++;
+  }
+
+  // Region 10
+  // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE10]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE10]]{{.+}})
+  // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK19-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8*
+  // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
+  // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]]
+  // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 20
+  // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]]
+
+  // CK19: call void [[CALL10:@.+]](i32* {{[^,]+}})
+  #pragma omp target map(tofrom:pa[20:60])
+  {
+    pa[50]++;
+  }
+
+  // Region 11
+  // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE11]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE11]]{{.+}})
+  // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK19-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8*
+  // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
+  // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]]
+  // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 0
+  // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]]
+
+  // CK19: call void [[CALL11:@.+]](i32* {{[^,]+}})
+  #pragma omp target map(alloc:pa[:60])
+  {
+    pa[50]++;
+  }
+
+  // Region 12
+  // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE12]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE12]]{{.+}})
+  // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK19-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8*
+  // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
+  // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]]
+  // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 15
+  // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]]
+
+  // CK19: call void [[CALL12:@.+]](i32* {{[^,]+}})
+  #pragma omp target map(to:pa[15])
+  {
+    pa[15]++;
+  }
+
+  // Region 13
+  // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE13]]{{.+}})
+  // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+  // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+
+  // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK19-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], i{{.+}}* [[S0]]
+  // CK19-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8*
+  // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
+  // CK19-DAG: [[CSVAL0]] = mul nuw i{{.+}} %{{.*}}, 4
+  // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]]
+  // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} %{{.*}}
+  // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]]
+
+  // CK19: call void [[CALL13:@.+]](i32* {{[^,]+}})
+  #pragma omp target map(alloc:pa[ii-23:ii])
+  {
+    pa[50]++;
+  }
+
+  // Region 14
+  // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE14]]{{.+}})
+  // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+  // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+
+  // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK19-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], i{{.+}}* [[S0]]
+  // CK19-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8*
+  // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
+  // CK19-DAG: [[CSVAL0]] = mul nuw i{{.+}} %{{.*}}, 4
+  // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]]
+  // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 0
+  // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]]
+
+  // CK19: call void [[CALL14:@.+]](i32* {{[^,]+}})
+  #pragma omp target map(to:pa[:ii])
+  {
+    pa[50]++;
+  }
+
+  // Region 15
+  // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE15]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE15]]{{.+}})
+  // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK19-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8*
+  // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
+  // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]]
+  // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} %{{.*}}
+  // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]]
+
+  // CK19: call void [[CALL15:@.+]](i32* {{[^,]+}})
+  #pragma omp target map(from:pa[ii+12])
+  {
+    pa[15]++;
+  }
+
+  // Map of a variable-size array.
+  int va[ii];
+
+  // Region 16
+  // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE16]]{{.+}})
+  // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+  // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+
+  // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK19-DAG: store i{{.+}} {{8|4}}, i{{.+}}* [[S0]]
+  // CK19-DAG: [[CBPVAL0]] = inttoptr i[[Z]] %{{.+}} to i8*
+  // CK19-DAG: [[CPVAL0]] = inttoptr i[[Z]] %{{.+}}to i8*
+
+  // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+  // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK19-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
+  // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]]
+  // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]]
+  // CK19-DAG: store i{{.+}} [[CSVAL1:%[^,]+]], i{{.+}}* [[S1]]
+  // CK19-DAG: [[CBPVAL1]] = bitcast i32* [[VAR1:%.+]] to i8*
+  // CK19-DAG: [[CPVAL1]] = bitcast i32* [[VAR1]] to i8*
+  // CK19-DAG: [[CSVAL1]] = mul nuw i{{.+}} %{{.*}}, 4
+
+  // CK19: call void [[CALL16:@.+]](i{{.+}} {{[^,]+}}, i32* {{[^,]+}})
+  #pragma omp target map(to:va)
+  {
+   va[50]++;
+  }
+
+  // Region 17
+  // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE17]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE17]]{{.+}})
+  // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK19-DAG: [[CBPVAL0]] = inttoptr i[[Z]] %{{.+}} to i8*
+  // CK19-DAG: [[CPVAL0]] = inttoptr i[[Z]] %{{.+}}to i8*
+
+  // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+  // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]]
+  // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]]
+  // CK19-DAG: [[CBPVAL1]] = bitcast i32* [[VAR1:%.+]] to i8*
+  // CK19-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8*
+  // CK19-DAG: [[SEC1]] = getelementptr {{.*}}i32* [[VAR1]], i{{.+}} 20
+
+  // CK19: call void [[CALL17:@.+]](i{{.+}} {{[^,]+}}, i32* {{[^,]+}})
+  #pragma omp target map(from:va[20:60])
+  {
+   va[50]++;
+  }
+
+  // Region 18
+  // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE18]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE18]]{{.+}})
+  // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK19-DAG: [[CBPVAL0]] = inttoptr i[[Z]] %{{.+}} to i8*
+  // CK19-DAG: [[CPVAL0]] = inttoptr i[[Z]] %{{.+}}to i8*
+
+  // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+  // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]]
+  // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]]
+  // CK19-DAG: [[CBPVAL1]] = bitcast i32* [[VAR1:%.+]] to i8*
+  // CK19-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8*
+  // CK19-DAG: [[SEC1]] = getelementptr {{.*}}i32* [[VAR1]], i{{.+}} 0
+
+  // CK19: call void [[CALL18:@.+]](i{{.+}} {{[^,]+}}, i32* {{[^,]+}})
+  #pragma omp target map(tofrom:va[:60])
+  {
+   va[50]++;
+  }
+
+  // Region 19
+  // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE19]]{{.+}})
+  // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+  // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+
+  // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK19-DAG: store i{{.+}} {{8|4}}, i{{.+}}* [[S0]]
+  // CK19-DAG: [[CBPVAL0]] = inttoptr i[[Z]] %{{.+}} to i8*
+  // CK19-DAG: [[CPVAL0]] = inttoptr i[[Z]] %{{.+}}to i8*
+
+  // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+  // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK19-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
+  // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]]
+  // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]]
+  // CK19-DAG: store i{{.+}} [[CSVAL1:%[^,]+]], i{{.+}}* [[S1]]
+  // CK19-DAG: [[CBPVAL1]] = bitcast i32* [[VAR1:%.+]] to i8*
+  // CK19-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8*
+  // CK19-DAG: [[CSVAL1]] = mul nuw i{{.+}} %{{.*}}, 4
+  // CK19-DAG: [[SEC1]] = getelementptr {{.*}}i32* [[VAR1]], i{{.+}} 0
+
+  // CK19: call void [[CALL19:@.+]](i{{.+}} {{[^,]+}}, i32* {{[^,]+}})
+  #pragma omp target map(alloc:va[:])
+  {
+   va[50]++;
+  }
+
+  // Region 20
+  // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE20]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE20]]{{.+}})
+  // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK19-DAG: [[CBPVAL0]] = inttoptr i[[Z]] %{{.+}} to i8*
+  // CK19-DAG: [[CPVAL0]] = inttoptr i[[Z]] %{{.+}}to i8*
+
+  // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+  // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]]
+  // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]]
+  // CK19-DAG: [[CBPVAL1]] = bitcast i32* [[VAR1:%.+]] to i8*
+  // CK19-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8*
+  // CK19-DAG: [[SEC1]] = getelementptr {{.*}}i32* [[VAR1]], i{{.+}} 15
+
+  // CK19: call void [[CALL20:@.+]](i{{.+}} {{[^,]+}}, i32* {{[^,]+}})
+  #pragma omp target map(to:va[15])
+  {
+   va[15]++;
+  }
+
+  // Region 21
+  // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE21]]{{.+}})
+  // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+  // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+
+  // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK19-DAG: store i{{.+}} {{8|4}}, i{{.+}}* [[S0]]
+  // CK19-DAG: [[CBPVAL0]] = inttoptr i[[Z]] %{{.+}} to i8*
+  // CK19-DAG: [[CPVAL0]] = inttoptr i[[Z]] %{{.+}}to i8*
+
+  // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+  // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK19-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
+  // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]]
+  // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]]
+  // CK19-DAG: store i{{.+}} [[CSVAL1:%[^,]+]], i{{.+}}* [[S1]]
+  // CK19-DAG: [[CBPVAL1]] = bitcast i32* [[VAR1:%.+]] to i8*
+  // CK19-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8*
+  // CK19-DAG: [[CSVAL1]] = mul nuw i{{.+}} %{{.*}}, 4
+  // CK19-DAG: [[SEC1]] = getelementptr {{.*}}i32* [[VAR1]], i{{.+}} %{{.+}}
+
+  // CK19: call void [[CALL21:@.+]](i{{.+}} {{[^,]+}}, i32* {{[^,]+}})
+  #pragma omp target map(tofrom:va[ii:ii+23])
+  {
+   va[50]++;
+  }
+
+  // Region 22
+  // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE22]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE22]]{{.+}})
+  // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK19-DAG: [[CBPVAL0]] = inttoptr i[[Z]] %{{.+}} to i8*
+  // CK19-DAG: [[CPVAL0]] = inttoptr i[[Z]] %{{.+}}to i8*
+
+  // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+  // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]]
+  // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]]
+  // CK19-DAG: [[CBPVAL1]] = bitcast i32* [[VAR1:%.+]] to i8*
+  // CK19-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8*
+  // CK19-DAG: [[SEC1]] = getelementptr {{.*}}i32* [[VAR1]], i{{.+}} %{{.+}}
+
+  // CK19: call void [[CALL22:@.+]](i{{.+}} {{[^,]+}}, i32* {{[^,]+}})
+  #pragma omp target map(tofrom:va[ii])
+  {
+   va[15]++;
+  }
+
+  // Always.
+  // Region 23
+  // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE23]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE23]]{{.+}})
+  // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK19-DAG: [[CBPVAL0]] = bitcast i32* [[VAR0:%.+]] to i8*
+  // CK19-DAG: [[CPVAL0]] = bitcast i32* [[VAR0]] to i8*
+
+  // CK19: call void [[CALL23:@.+]](i32* {{[^,]+}})
+  #pragma omp target map(always, tofrom: a)
+  {
+   a++;
+  }
+}
+
+// CK19: define {{.+}}[[CALL00]]
+// CK19: define {{.+}}[[CALL01]]
+// CK19: define {{.+}}[[CALL02]]
+// CK19: define {{.+}}[[CALL03]]
+// CK19: define {{.+}}[[CALL04]]
+// CK19: define {{.+}}[[CALL05]]
+// CK19: define {{.+}}[[CALL06]]
+// CK19: define {{.+}}[[CALL07]]
+// CK19: define {{.+}}[[CALL08]]
+// CK19: define {{.+}}[[CALL09]]
+// CK19: define {{.+}}[[CALL10]]
+// CK19: define {{.+}}[[CALL11]]
+// CK19: define {{.+}}[[CALL12]]
+// CK19: define {{.+}}[[CALL13]]
+// CK19: define {{.+}}[[CALL14]]
+// CK19: define {{.+}}[[CALL15]]
+// CK19: define {{.+}}[[CALL16]]
+// CK19: define {{.+}}[[CALL17]]
+// CK19: define {{.+}}[[CALL18]]
+// CK19: define {{.+}}[[CALL19]]
+// CK19: define {{.+}}[[CALL20]]
+// CK19: define {{.+}}[[CALL21]]
+// CK19: define {{.+}}[[CALL22]]
+// CK19: define {{.+}}[[CALL23]]
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK20 -verify -fopenmp -omptargets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK20 --check-prefix CK20-64
+// RUN: %clang_cc1 -DCK20 -fopenmp -omptargets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -omptargets=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 CK20 --check-prefix CK20-64
+// RUN: %clang_cc1 -DCK20 -verify -fopenmp -omptargets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK20 --check-prefix CK20-32
+// RUN: %clang_cc1 -DCK20 -fopenmp -omptargets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -omptargets=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 CK20 --check-prefix CK20-32
+#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: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20]
+// CK20: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 1]
+
+// CK20: [[SIZE02:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
+// CK20: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 2]
+
+// CK20: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 12]
+// CK20: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 2]
+
+// CK20-LABEL: explicit_maps_references_and_function_args
+void explicit_maps_references_and_function_args (int a, float b, int (&c)[10], float *d){
+
+  int &aa = a;
+  float &bb = b;
+  int (&cc)[10] = c;
+  float *&dd = d;
+
+  // Region 00
+  // CK20-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+  // CK20-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK20-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK20-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK20-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK20-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK20-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK20-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8*
+  // CK20-DAG: [[CPVAL0]] = bitcast i32* [[RVAR00:%.+]] to i8*
+  // CK20-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]]
+  // CK20-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]]
+
+  // CK20: call void [[CALL00:@.+]](i32* {{[^,]+}})
+  #pragma omp target map(to:aa)
+  {
+    aa += 1;
+  }
+
+  // Region 01
+  // CK20-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
+  // CK20-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK20-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK20-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK20-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK20-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK20-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK20-DAG: [[CBPVAL0]] = bitcast [10 x i32]* [[RVAR0:%.+]] to i8*
+  // CK20-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
+  // CK20-DAG: [[SEC0]] = getelementptr {{.*}}[10 x i32]* [[RVAR00:%.+]], i{{.+}} 0, i{{.+}} 0
+  // CK20-DAG: [[RVAR0]] = load [10 x i32]*, [10 x i32]** [[VAR0:%[^,]+]]
+  // CK20-DAG: [[RVAR00]] = load [10 x i32]*, [10 x i32]** [[VAR0]]
+
+  // CK20: call void [[CALL01:@.+]]([10 x i32]* {{[^,]+}})
+  #pragma omp target map(to:cc[:5])
+  {
+    cc[3] += 1;
+  }
+
+  // Region 02
+  // CK20-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}})
+  // CK20-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK20-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK20-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK20-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK20-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK20-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK20-DAG: [[CBPVAL0]] = bitcast float* [[VAR0:%.+]] to i8*
+  // CK20-DAG: [[CPVAL0]] = bitcast float* [[VAR0]] to i8*
+
+  // CK20: call void [[CALL02:@.+]](float* {{[^,]+}})
+  #pragma omp target map(from:b)
+  {
+    b += 1.0f;
+  }
+
+  // Region 03
+  // CK20-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE03]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}})
+  // CK20-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK20-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK20-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK20-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK20-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK20-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK20-DAG: [[CBPVAL0]] = bitcast float* [[RVAR0:%.+]] to i8*
+  // CK20-DAG: [[CPVAL0]] = bitcast float* [[SEC0:%.+]] to i8*
+  // CK20-DAG: [[RVAR0]] = load float*, float** [[VAR0:%[^,]+]]
+  // CK20-DAG: [[SEC0]] = getelementptr {{.*}}float* [[RVAR00:%.+]], i{{.+}} 2
+  // CK20-DAG: [[RVAR00]] = load float*, float** [[VAR0]]
+
+  // CK20: call void [[CALL03:@.+]](float* {{[^,]+}})
+  #pragma omp target map(from:d[2:3])
+  {
+    d[2] += 1.0f;
+  }
+}
+
+// CK20: define {{.+}}[[CALL00]]
+// CK20: define {{.+}}[[CALL01]]
+// CK20: define {{.+}}[[CALL02]]
+// CK20: define {{.+}}[[CALL03]]
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK21 -verify -fopenmp -omptargets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK21 --check-prefix CK21-64
+// RUN: %clang_cc1 -DCK21 -fopenmp -omptargets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -omptargets=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 CK21 --check-prefix CK21-64
+// RUN: %clang_cc1 -DCK21 -verify -fopenmp -omptargets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK21 --check-prefix CK21-32
+// RUN: %clang_cc1 -DCK21 -fopenmp -omptargets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -omptargets=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 CK21 --check-prefix CK21-32
+#ifdef CK21
+// 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: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 492]
+// CK21: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// 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: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 492]
+// CK21: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 2]
+
+// CK21: [[SIZE04:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
+// CK21: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 2]
+
+// 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-LABEL: explicit_maps_template_args_and_members
+
+template <int X, typename T>
+struct CC {
+  T A;
+  int A2;
+  float *B;
+
+  int foo(T arg) {
+    float la[X];
+    T *lb;
+
+    // Region 00
+    // CK21-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+    // CK21-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+    // CK21-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+    // CK21-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+    // CK21-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+    // CK21-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+    // CK21-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+    // CK21-DAG: [[CBPVAL0]] = bitcast [[ST]]* [[VAR0:%.+]] to i8*
+    // CK21-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
+    // CK21-DAG: [[SEC0]] = getelementptr {{.*}}[[ST]]* [[VAR0:%.+]], i{{.+}} 0, i{{.+}} 0
+
+    // CK21: call void [[CALL00:@.+]]([[ST]]* {{[^,]+}})
+    #pragma omp target map(A)
+    {
+      A += 1;
+    }
+    
+    // Region 01
+    // CK21-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
+    // CK21-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+    // CK21-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+    // CK21-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+    // CK21-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+    // CK21-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+    // CK21-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+    // CK21-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8*
+    // CK21-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
+    // CK21-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]]
+    // CK21-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 0
+    // CK21-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]]
+
+    // CK21: call void [[CALL01:@.+]](i32* {{[^,]+}})
+    #pragma omp target map(lb[:X])
+    {
+      lb[4] += 1;
+    }
+    
+    // Region 02
+    // CK21-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE02]]{{.+}})
+    // CK21-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+    // CK21-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+    // CK21-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+    // CK21-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+    // CK21-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+    // CK21-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+    // CK21-DAG: [[CBPVAL0]] = bitcast [[ST]]* [[VAR0:%.+]] to i8*
+    // CK21-DAG: [[CPVAL0]] = bitcast float** [[SEC0:%.+]] to i8*
+    // CK21-DAG: [[SEC0]] = getelementptr {{.*}}[[ST]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2
+    
+    // CK21-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+    // CK21-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+    // CK21-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]]
+    // CK21-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]]
+    // CK21-DAG: [[CBPVAL1]] = bitcast float** [[SEC0]] to i8*
+    // CK21-DAG: [[CPVAL1]] = bitcast float* [[SEC1:%.+]] to i8*
+    // CK21-DAG: [[SEC1]] = getelementptr {{.*}}float* [[RVAR1:%[^,]+]], i{{.+}} 123
+    // CK21-DAG: [[RVAR1]] = load float*, float** [[SEC1_:%[^,]+]]
+    // CK21-DAG: [[SEC1_]] = getelementptr {{.*}}[[ST]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2
+
+    // CK21: call void [[CALL02:@.+]]([[ST]]* {{[^,]+}})
+    #pragma omp target map(from:B[X:X+2])
+    {
+      B[2] += 1.0f;
+    }
+    
+    // Region 03
+    // CK21-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE03]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}})
+    // CK21-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+    // CK21-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+    // CK21-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+    // CK21-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+    // CK21-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+    // CK21-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+    // CK21-DAG: [[CBPVAL0]] = bitcast [123 x float]* [[VAR0:%.+]] to i8*
+    // CK21-DAG: [[CPVAL0]] = bitcast [123 x float]* [[VAR0]] to i8*
+
+    // CK21: call void [[CALL03:@.+]]([123 x float]* {{[^,]+}})
+    #pragma omp target map(from:la)
+    {
+      la[3] += 1.0f;
+    }
+    
+    // Region 04
+    // CK21-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE04]]{{.+}})
+    // CK21-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+    // CK21-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+    // CK21-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+    // CK21-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+    // CK21-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+    // CK21-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+    // CK21-DAG: [[CBPVAL0]] = bitcast i32* [[VAR0:%.+]] to i8*
+    // CK21-DAG: [[CPVAL0]] = bitcast i32* [[VAR0]] to i8*
+
+    // CK21: call void [[CALL04:@.+]](i32* {{[^,]+}})
+    #pragma omp target map(from:arg)
+    {
+      arg +=1;
+    }
+    
+    // Make sure the extra flag is passed to the second map.
+    // Region 05
+    // CK21-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE05]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE05]]{{.+}})
+    // CK21-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+    // CK21-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+    // CK21-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+    // CK21-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+    // CK21-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+    // CK21-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+    // CK21-DAG: [[CBPVAL0]] = bitcast [[ST]]* [[VAR0:%.+]] to i8*
+    // CK21-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
+    // CK21-DAG: [[SEC0]] = getelementptr {{.*}}[[ST]]* [[VAR0]], i{{.+}} 0, i{{.+}} 0
+
+    // CK21-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+    // CK21-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+    // CK21-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]]
+    // CK21-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]]
+    // CK21-DAG: [[CBPVAL1]] = bitcast [[ST]]* [[VAR1:%.+]] to i8*
+    // CK21-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8*
+    // CK21-DAG: [[SEC1]] = getelementptr {{.*}}[[ST]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1
+
+    // CK21: call void [[CALL05:@.+]]([[ST]]* {{[^,]+}})
+    #pragma omp target map(A, A2)
+    {
+      A += 1;
+      A2 += 1;
+    }
+    return A;
+  }
+};
+
+int explicit_maps_template_args_and_members(int a){
+  CC<123,int> c;
+  return c.foo(a);
+}
+
+// CK21: define {{.+}}[[CALL00]]
+// CK21: define {{.+}}[[CALL01]]
+// CK21: define {{.+}}[[CALL02]]
+// CK21: define {{.+}}[[CALL03]]
+// CK21: define {{.+}}[[CALL04]]
+// CK21: define {{.+}}[[CALL05]]
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK22 -verify -fopenmp -omptargets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK22 --check-prefix CK22-64
+// RUN: %clang_cc1 -DCK22 -fopenmp -omptargets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -omptargets=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 CK22 --check-prefix CK22-64
+// RUN: %clang_cc1 -DCK22 -verify -fopenmp -omptargets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK22 --check-prefix CK22-32
+// RUN: %clang_cc1 -DCK22 -fopenmp -omptargets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -omptargets=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 CK22 --check-prefix CK22-32
+#ifdef CK22
+
+// CK22-DAG: [[ST:%.+]] = type { float }
+// 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: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400]
+// CK22: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// CK22: [[SIZE02:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{8|4}}]
+// CK22: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// CK22: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 16]
+// CK22: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// CK22: [[SIZE04:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20]
+// CK22: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// CK22: [[SIZE05:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
+// CK22: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// CK22: [[SIZE06:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400]
+// CK22: [[MTYPE06:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// CK22: [[SIZE07:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{8|4}}]
+// CK22: [[MTYPE07:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// CK22: [[SIZE08:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 16]
+// CK22: [[MTYPE08:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// CK22: [[SIZE09:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20]
+// CK22: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// CK22: [[SIZE10:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
+// CK22: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// CK22: [[SIZE11:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400]
+// CK22: [[MTYPE11:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// CK22: [[SIZE12:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{8|4}}]
+// CK22: [[MTYPE12:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// CK22: [[SIZE13:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 16]
+// CK22: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// CK22: [[SIZE14:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20]
+// CK22: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+int a;
+int c[100];
+int *d;
+
+struct ST {
+  float fa;
+};
+
+ST sa ;
+ST sc[100];
+ST *sd;
+
+template<typename T>
+struct STT {
+  T fa;
+};
+
+STT<int> sta ;
+STT<int> stc[100];
+STT<int> *std;
+
+// CK22-LABEL: explicit_maps_globals
+int explicit_maps_globals(void){
+  // Region 00
+  // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+  // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK22-DAG: store i8* bitcast (i32* @a to i8*), i8** [[BP0]]
+  // CK22-DAG: store i8* bitcast (i32* @a to i8*), i8** [[P0]]
+
+  // CK22: call void [[CALL00:@.+]](i32* {{[^,]+}})
+  #pragma omp target map(a)
+  { a+=1; }
+  
+  // Region 01
+  // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
+  // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK22-DAG: store i8* bitcast ([100 x i32]* @c to i8*), i8** [[BP0]]
+  // CK22-DAG: store i8* bitcast ([100 x i32]* @c to i8*), i8** [[P0]]
+
+  // CK22: call void [[CALL01:@.+]]([100 x i32]* {{[^,]+}})
+  #pragma omp target map(c)
+  { c[3]+=1; }
+  
+  // Region 02
+  // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}})
+  // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK22-DAG: store i8* bitcast (i32** @d to i8*), i8** [[BP0]]
+  // CK22-DAG: store i8* bitcast (i32** @d to i8*), i8** [[P0]]
+
+  // CK22: call void [[CALL02:@.+]](i32** {{[^,]+}})
+  #pragma omp target map(d)
+  { d[3]+=1; }
+    
+  // Region 03
+  // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE03]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}})
+  // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK22-DAG: store i8* bitcast ([100 x i32]* @c to i8*), i8** [[BP0]]
+  // CK22-DAG: store i8* bitcast (i32* getelementptr inbounds ([100 x i32], [100 x i32]* @c, i{{.+}} 0, i{{.+}} 1) to i8*), i8** [[P0]]
+
+  // CK22: call void [[CALL03:@.+]]([100 x i32]* {{[^,]+}})
+  #pragma omp target map(c[1:4])
+  { c[3]+=1; }
+  
+  // Region 04
+  // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE04]]{{.+}})
+  // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK22-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK22-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK22-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8*
+  // CK22-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
+  // CK22-DAG: [[RVAR0]] = load i32*, i32** @d
+  // CK22-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 2
+  // CK22-DAG: [[RVAR00]] = load i32*, i32** @d
+
+  // CK22: call void [[CALL04:@.+]](i32* {{[^,]+}})
+  #pragma omp target map(d[2:5])
+  { d[3]+=1; }
+  
+  // Region 05
+  // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE05]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE05]]{{.+}})
+  // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK22-DAG: store i8* bitcast ([[ST]]* @sa to i8*), i8** [[BP0]]
+  // CK22-DAG: store i8* bitcast ([[ST]]* @sa to i8*), i8** [[P0]]
+
+  // CK22: call void [[CALL05:@.+]]([[ST]]* {{[^,]+}})
+  #pragma omp target map(sa)
+  { sa.fa+=1; }
+  
+  // Region 06
+  // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE06]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE06]]{{.+}})
+  // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK22-DAG: store i8* bitcast ([100 x [[ST]]]* @sc to i8*), i8** [[BP0]]
+  // CK22-DAG: store i8* bitcast ([100 x [[ST]]]* @sc to i8*), i8** [[P0]]
+
+  // CK22: call void [[CALL06:@.+]]([100 x [[ST]]]* {{[^,]+}})
+  #pragma omp target map(sc)
+  { sc[3].fa+=1; }
+  
+  // Region 07
+  // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE07]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE07]]{{.+}})
+  // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK22-DAG: store i8* bitcast ([[ST]]** @sd to i8*), i8** [[BP0]]
+  // CK22-DAG: store i8* bitcast ([[ST]]** @sd to i8*), i8** [[P0]]
+
+  // CK22: call void [[CALL07:@.+]]([[ST]]** {{[^,]+}})
+  #pragma omp target map(sd)
+  { sd[3].fa+=1; }
+  
+  // Region 08
+  // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE08]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE08]]{{.+}})
+  // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK22-DAG: store i8* bitcast ([100 x [[ST]]]* @sc to i8*), i8** [[BP0]]
+  // CK22-DAG: store i8* bitcast ([[ST]]* getelementptr inbounds ([100 x [[ST]]], [100 x [[ST]]]* @sc, i{{.+}} 0, i{{.+}} 1) to i8*), i8** [[P0]]
+
+  // CK22: call void [[CALL08:@.+]]([100 x [[ST]]]* {{[^,]+}})
+  #pragma omp target map(sc[1:4])
+  { sc[3].fa+=1; }
+  
+  // Region 09
+  // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE09]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE09]]{{.+}})
+  // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK22-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK22-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK22-DAG: [[CBPVAL0]] = bitcast [[ST]]* [[RVAR0:%.+]] to i8*
+  // CK22-DAG: [[CPVAL0]] = bitcast [[ST]]* [[SEC0:%.+]] to i8*
+  // CK22-DAG: [[RVAR0]] = load [[ST]]*, [[ST]]** @sd
+  // CK22-DAG: [[SEC0]] = getelementptr {{.*}}[[ST]]* [[RVAR00:%.+]], i{{.+}} 2
+  // CK22-DAG: [[RVAR00]] = load [[ST]]*, [[ST]]** @sd
+
+  // CK22: call void [[CALL09:@.+]]([[ST]]* {{[^,]+}})
+  #pragma omp target map(sd[2:5])
+  { sd[3].fa+=1; }
+  
+  // Region 10
+  // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE10]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE10]]{{.+}})
+  // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK22-DAG: store i8* bitcast ([[STT]]* @sta to i8*), i8** [[BP0]]
+  // CK22-DAG: store i8* bitcast ([[STT]]* @sta to i8*), i8** [[P0]]
+
+  // CK22: call void [[CALL10:@.+]]([[STT]]* {{[^,]+}})
+  #pragma omp target map(sta)
+  { sta.fa+=1; }
+  
+  // Region 11
+  // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE11]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE11]]{{.+}})
+  // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK22-DAG: store i8* bitcast ([100 x [[STT]]]* @stc to i8*), i8** [[BP0]]
+  // CK22-DAG: store i8* bitcast ([100 x [[STT]]]* @stc to i8*), i8** [[P0]]
+
+  // CK22: call void [[CALL11:@.+]]([100 x [[STT]]]* {{[^,]+}})
+  #pragma omp target map(stc)
+  { stc[3].fa+=1; }
+  
+  // Region 12
+  // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE12]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE12]]{{.+}})
+  // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK22-DAG: store i8* bitcast ([[STT]]** @std to i8*), i8** [[BP0]]
+  // CK22-DAG: store i8* bitcast ([[STT]]** @std to i8*), i8** [[P0]]
+
+  // CK22: call void [[CALL12:@.+]]([[STT]]** {{[^,]+}})
+  #pragma omp target map(std)
+  { std[3].fa+=1; }
+  
+  // Region 13
+  // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE13]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE13]]{{.+}})
+  // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK22-DAG: store i8* bitcast ([100 x [[STT]]]* @stc to i8*), i8** [[BP0]]
+  // CK22-DAG: store i8* bitcast ([[STT]]* getelementptr inbounds ([100 x [[STT]]], [100 x [[STT]]]* @stc, i{{.+}} 0, i{{.+}} 1) to i8*), i8** [[P0]]
+
+  // CK22: call void [[CALL13:@.+]]([100 x [[STT]]]* {{[^,]+}})
+  #pragma omp target map(stc[1:4])
+  { stc[3].fa+=1; }
+  
+  // Region 14
+  // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE14]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE14]]{{.+}})
+  // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK22-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK22-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK22-DAG: [[CBPVAL0]] = bitcast [[STT]]* [[RVAR0:%.+]] to i8*
+  // CK22-DAG: [[CPVAL0]] = bitcast [[STT]]* [[SEC0:%.+]] to i8*
+  // CK22-DAG: [[RVAR0]] = load [[STT]]*, [[STT]]** @std
+  // CK22-DAG: [[SEC0]] = getelementptr {{.*}}[[STT]]* [[RVAR00:%.+]], i{{.+}} 2
+  // CK22-DAG: [[RVAR00]] = load [[STT]]*, [[STT]]** @std
+
+  // CK22: call void [[CALL14:@.+]]([[STT]]* {{[^,]+}})
+  #pragma omp target map(std[2:5])
+  { std[3].fa+=1; }
+
+  return 0;
+}
+// CK22: define {{.+}}[[CALL00]]
+// CK22: define {{.+}}[[CALL01]]
+// CK22: define {{.+}}[[CALL02]]
+// CK22: define {{.+}}[[CALL03]]
+// CK22: define {{.+}}[[CALL04]]
+// CK22: define {{.+}}[[CALL05]]
+// CK22: define {{.+}}[[CALL06]]
+// CK22: define {{.+}}[[CALL07]]
+// CK22: define {{.+}}[[CALL08]]
+// CK22: define {{.+}}[[CALL09]]
+// CK22: define {{.+}}[[CALL10]]
+// CK22: define {{.+}}[[CALL11]]
+// CK22: define {{.+}}[[CALL12]]
+// CK22: define {{.+}}[[CALL13]]
+// CK22: define {{.+}}[[CALL14]]
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -std=c++11 -DCK23 -verify -fopenmp -omptargets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK23 --check-prefix CK23-64
+// RUN: %clang_cc1 -std=c++11 -DCK23 -fopenmp -omptargets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -std=c++11 -fopenmp -omptargets=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 CK23 --check-prefix CK23-64
+// RUN: %clang_cc1 -std=c++11 -DCK23 -verify -fopenmp -omptargets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK23 --check-prefix CK23-32
+// RUN: %clang_cc1 -std=c++11 -DCK23 -fopenmp -omptargets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -std=c++11 -fopenmp -omptargets=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 CK23 --check-prefix CK23-32
+#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: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
+// CK23: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// CK23: [[SIZE02:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400]
+// CK23: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// CK23: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{8|4}}]
+// CK23: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// CK23: [[SIZE04:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 16]
+// CK23: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// CK23: [[SIZE05:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 16]
+// CK23: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// CK23-LABEL: explicit_maps_inside_captured
+int explicit_maps_inside_captured(int a){
+  float b;
+  float c[100];
+  float *d;
+  
+  // CK23: call void @{{.*}}explicit_maps_inside_captured{{.*}}([[SA:%.+]]* {{.*}}) 
+  // CK23: define {{.*}}explicit_maps_inside_captured{{.*}}
+  [&](void){
+    // Region 00
+    // CK23-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+    // CK23-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+    // CK23-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+    // CK23-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+    // CK23-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+    // CK23-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+    // CK23-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+    // CK23-DAG: [[CBPVAL0]] = bitcast i32* [[VAR0:%.+]] to i8*
+    // CK23-DAG: [[CPVAL0]] = bitcast i32* [[VAR00:%.+]] to i8*
+    // CK23-DAG: [[VAR0]] = load i32*, i32** [[CAP0:%[^,]+]]
+    // CK23-DAG: [[CAP0]] = getelementptr inbounds [[SA]], [[SA]]
+    // CK23-DAG: [[VAR00]] = load i32*, i32** [[CAP00:%[^,]+]]
+    // CK23-DAG: [[CAP00]] = getelementptr inbounds [[SA]], [[SA]]
+
+    // CK23: call void [[CALL00:@.+]](i32* {{[^,]+}})
+    #pragma omp target map(a)
+      { a+=1; }
+    // Region 01
+    // CK23-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
+    // CK23-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+    // CK23-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+    // CK23-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+    // CK23-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+    // CK23-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+    // CK23-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+    // CK23-DAG: [[CBPVAL0]] = bitcast float* [[VAR0:%.+]] to i8*
+    // CK23-DAG: [[CPVAL0]] = bitcast float* [[VAR00:%.+]] to i8*
+    // CK23-DAG: [[VAR0]] = load float*, float** [[CAP0:%[^,]+]]
+    // CK23-DAG: [[CAP0]] = getelementptr inbounds [[SA]], [[SA]]
+    // CK23-DAG: [[VAR00]] = load float*, float** [[CAP00:%[^,]+]]
+    // CK23-DAG: [[CAP00]] = getelementptr inbounds [[SA]], [[SA]]
+
+    // CK23: call void [[CALL01:@.+]](float* {{[^,]+}})
+    #pragma omp target map(b)
+      { b+=1; }
+    // Region 02
+    // CK23-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}})
+    // CK23-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+    // CK23-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+    // CK23-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+    // CK23-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+    // CK23-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+    // CK23-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+    // CK23-DAG: [[CBPVAL0]] = bitcast [100 x float]* [[VAR0:%.+]] to i8*
+    // CK23-DAG: [[CPVAL0]] = bitcast [100 x float]* [[VAR00:%.+]] to i8*
+    // CK23-DAG: [[VAR0]] = load [100 x float]*, [100 x float]** [[CAP0:%[^,]+]]
+    // CK23-DAG: [[CAP0]] = getelementptr inbounds [[SA]], [[SA]]
+    // CK23-DAG: [[VAR00]] = load [100 x float]*, [100 x float]** [[CAP00:%[^,]+]]
+    // CK23-DAG: [[CAP00]] = getelementptr inbounds [[SA]], [[SA]]
+
+    // CK23: call void [[CALL02:@.+]]([100 x float]* {{[^,]+}})
+    #pragma omp target map(c)
+      { c[3]+=1; }
+      
+    // Region 03
+    // CK23-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE03]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}})
+    // CK23-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+    // CK23-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+    // CK23-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+    // CK23-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+    // CK23-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+    // CK23-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+    // CK23-DAG: [[CBPVAL0]] = bitcast float** [[VAR0:%.+]] to i8*
+    // CK23-DAG: [[CPVAL0]] = bitcast float** [[VAR00:%.+]] to i8*
+    // CK23-DAG: [[VAR0]] = load float**, float*** [[CAP0:%[^,]+]]
+    // CK23-DAG: [[CAP0]] = getelementptr inbounds [[SA]], [[SA]]
+    // CK23-DAG: [[VAR00]] = load float**, float*** [[CAP00:%[^,]+]]
+    // CK23-DAG: [[CAP00]] = getelementptr inbounds [[SA]], [[SA]]
+
+    // CK23: call void [[CALL03:@.+]](float** {{[^,]+}})
+    #pragma omp target map(d)
+      { d[3]+=1; }
+    // Region 04
+    // CK23-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE04]]{{.+}})
+    // CK23-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+    // CK23-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+    // CK23-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+    // CK23-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+    // CK23-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+    // CK23-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+    // CK23-DAG: [[CBPVAL0]] = bitcast [100 x float]* [[VAR0:%.+]] to i8*
+    // CK23-DAG: [[CPVAL0]] = bitcast float* [[SEC0:%.+]] to i8*
+    // CK23-DAG: [[SEC0]] = getelementptr {{.*}}[100 x float]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 2
+    // CK23-DAG: [[VAR0]] = load [100 x float]*, [100 x float]** [[CAP0:%[^,]+]]
+    // CK23-DAG: [[CAP0]] = getelementptr inbounds [[SA]], [[SA]]
+    // CK23-DAG: [[VAR00]] = load [100 x float]*, [100 x float]** [[CAP00:%[^,]+]]
+    // CK23-DAG: [[CAP00]] = getelementptr inbounds [[SA]], [[SA]]
+    
+    // CK23: call void [[CALL04:@.+]]([100 x float]* {{[^,]+}})
+    #pragma omp target map(c[2:4])
+      { c[3]+=1; }
+      
+    // Region 05
+    // CK23-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE05]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE05]]{{.+}})
+    // CK23-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+    // CK23-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+    // CK23-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+    // CK23-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+    // CK23-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+    // CK23-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+    // CK23-DAG: [[CBPVAL0]] = bitcast float* [[RVAR0:%.+]] to i8*
+    // CK23-DAG: [[CPVAL0]] = bitcast float* [[SEC0:%.+]] to i8*
+    // CK23-DAG: [[RVAR0]] = load float*, float** [[VAR0:%[^,]+]]
+    // CK23-DAG: [[SEC0]] = getelementptr {{.*}}float* [[RVAR00:%.+]], i{{.+}} 2
+    // CK23-DAG: [[RVAR00]] = load float*, float** [[VAR00:%[^,]+]]
+    // CK23-DAG: [[VAR0]] = load float**, float*** [[CAP0:%[^,]+]]
+    // CK23-DAG: [[CAP0]] = getelementptr inbounds [[SA]], [[SA]]
+    // CK23-DAG: [[VAR00]] = load float**, float*** [[CAP00:%[^,]+]]
+    // CK23-DAG: [[CAP00]] = getelementptr inbounds [[SA]], [[SA]]
+        
+    // CK23: call void [[CALL05:@.+]](float* {{[^,]+}})
+    #pragma omp target map(d[2:4])
+      { d[3]+=1; }
+  }();
+  return b;
+}
+
+// CK23: define {{.+}}[[CALL00]]
+// CK23: define {{.+}}[[CALL01]]
+// CK23: define {{.+}}[[CALL02]]
+// CK23: define {{.+}}[[CALL03]]
+// CK23: define {{.+}}[[CALL04]]
+// CK23: define {{.+}}[[CALL05]]
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK24 -verify -fopenmp -omptargets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK24 --check-prefix CK24-64
+// RUN: %clang_cc1 -DCK24 -fopenmp -omptargets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -omptargets=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 CK24 --check-prefix CK24-64
+// RUN: %clang_cc1 -DCK24 -verify -fopenmp -omptargets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK24 --check-prefix CK24-32
+// RUN: %clang_cc1 -DCK24 -fopenmp -omptargets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -omptargets=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 CK24 --check-prefix CK24-32
+#ifdef CK24
+
+// CK24-DAG: [[SC:%.+]] = type { i32, [[SB:%.+]], [[SB:%.+]]*, [10 x i32] }
+// CK24-DAG: [[SB]] = type { i32, [[SA:%.+]], [10 x [[SA:%.+]]], [10 x [[SA:%.+]]*], [[SA:%.+]]* }
+// CK24-DAG: [[SA]] = type { i32, [[SA]]*, [10 x i32] }
+
+struct SA{
+  int a;
+  struct SA *p;
+  int b[10];
+};
+struct SB{
+  int a;
+  struct SA s;
+  struct SA sa[10];
+  struct SA *sp[10];
+  struct SA *p;
+};
+struct SC{
+  int a;
+  struct SB s;
+  struct SB *p;
+  int b[10];
+};
+
+// 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: [[SIZE02:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{56|48}}]
+// CK24: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// CK24: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
+// CK24: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// CK24: [[SIZE04:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20]
+// CK24: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// 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: [[SIZE06:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
+// CK24: [[MTYPE06:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// 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: [[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: [[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: [[SIZE10:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 8]
+// CK24: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// 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: [[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: [[SIZE13:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
+// CK24: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// CK24: [[SIZE14:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{56|48}}]
+// CK24: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// CK24: [[SIZE15:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
+// CK24: [[MTYPE15:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// CK24: [[SIZE16:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20]
+// CK24: [[MTYPE16:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// 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: [[SIZE18:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
+// CK24: [[MTYPE18:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// 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: [[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: [[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: [[SIZE22:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{8|4}}]
+// CK24: [[MTYPE22:@.+]] = private {{.*}}constant [1 x i32] [i32 3]
+
+// 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: [[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-LABEL: explicit_maps_struct_fields
+int explicit_maps_struct_fields(int a){
+  SC s;
+  SC *p;
+
+// Region 01
+// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
+// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8*
+// CK24-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
+// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 0
+
+// CK24: call void [[CALL01:@.+]]([[SC]]* {{[^,]+}})
+#pragma omp target map(s.a)
+  { s.a++; }
+  
+// Region 02
+// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}})
+// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8*
+// CK24-DAG: [[CPVAL0]] = bitcast [[SA]]* [[SEC0:%.+]] to i8*
+// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SB]]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1
+
+// CK24: call void [[CALL02:@.+]]([[SC]]* {{[^,]+}})
+#pragma omp target map(s.s.s)
+  { s.a++; }
+  
+// Region 03
+// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE03]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}})
+// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8*
+// CK24-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
+// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SA]]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SB]]* [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: [[SEC000]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1
+
+// CK24: call void [[CALL03:@.+]]([[SC]]* {{[^,]+}})
+#pragma omp target map(s.s.s.a)
+  { s.a++; }
+
+// Region 04
+// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE04]]{{.+}})
+// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8*
+// CK24-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
+// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[10 x i32]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 3
+
+// CK24: call void [[CALL04:@.+]]([[SC]]* {{[^,]+}})
+#pragma omp target map(s.b[:5])
+  { s.a++; }
+  
+// Region 05
+// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE05]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE05]]{{.+}})
+// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8*
+// CK24-DAG: [[CPVAL0]] = bitcast [[SB]]** [[SEC0:%.+]] to i8*
+// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2
+
+// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]]
+// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]]
+// CK24-DAG: [[CBPVAL1]] = bitcast [[SB]]** [[SEC0]] to i8*
+// CK24-DAG: [[CPVAL1]] = bitcast [[SB]]* [[SEC1:%.+]] to i8*
+// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SB]]* [[SEC11:%[^,]+]], i{{.+}} 0
+// CK24-DAG: [[SEC11]] = load [[SB]]*, [[SB]]** [[SEC111:%[^,]+]],
+// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2
+
+// CK24: call void [[CALL05:@.+]]([[SC]]* {{[^,]+}})  
+#pragma omp target map(s.p[:5])
+  { s.a++; }
+
+// Region 06
+// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE06]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE06]]{{.+}})
+// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8*
+// CK24-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
+// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SA]]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[10 x [[SA]]]* [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 3
+// CK24-DAG: [[SEC000]] = getelementptr {{.*}}[[SB]]* [[SEC0000:%[^,]+]], i{{.+}} 0, i{{.+}} 2
+// CK24-DAG: [[SEC0000]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1
+
+// CK24: call void [[CALL06:@.+]]([[SC]]* {{[^,]+}})
+#pragma omp target map(s.s.sa[3].a)
+  { s.a++; }
+  
+// Region 07
+// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE07]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE07]]{{.+}})
+// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8*
+// CK24-DAG: [[CPVAL0]] = bitcast [[SA]]** [[SEC0:%.+]] to i8*
+// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[10 x [[SA]]*]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 3
+// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SB]]* [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 3
+// CK24-DAG: [[SEC000]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1
+
+// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]]
+// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]]
+// CK24-DAG: [[CBPVAL1]] = bitcast [[SA]]** [[SEC0]] to i8*
+// CK24-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8*
+// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SA]]* [[SEC11:%[^,]+]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[SEC11]] = load [[SA]]*, [[SA]]** [[SEC111:%[^,]+]],
+// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[10 x [[SA]]*]* [[SEC1111:%[^,]+]], i{{.+}} 0, i{{.+}} 3
+// CK24-DAG: [[SEC1111]] = getelementptr {{.*}}[[SB]]* [[SEC11111:%[^,]+]], i{{.+}} 0, i{{.+}} 3
+// CK24-DAG: [[SEC11111]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1
+
+// CK24: call void [[CALL07:@.+]]([[SC]]* {{[^,]+}}) 
+#pragma omp target map(s.s.sp[3]->a)
+  { s.a++; }
+
+// Region 08
+// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE08]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE08]]{{.+}})
+// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8*
+// CK24-DAG: [[CPVAL0]] = bitcast [[SB]]** [[SEC0:%.+]] to i8*
+// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2
+
+// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]]
+// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]]
+// CK24-DAG: [[CBPVAL1]] = bitcast [[SB]]** [[SEC0]] to i8*
+// CK24-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8*
+// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SB]]* [[SEC11:%[^,]+]], i{{.+}} 0
+// CK24-DAG: [[SEC11]] = load [[SB]]*, [[SB]]** [[SEC111:%[^,]+]],
+// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2
+
+// CK24: call void [[CALL08:@.+]]([[SC]]* {{[^,]+}}) 
+#pragma omp target map(s.p->a)
+  { s.a++; }
+  
+// Region 09
+// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE09]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE09]]{{.+}})
+// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8*
+// CK24-DAG: [[CPVAL0]] = bitcast [[SA]]** [[SEC0:%.+]] to i8*
+// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SB]]* [[SEC00:[^,]+]], i{{.+}} 0, i{{.+}} 4
+// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1
+
+// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]]
+// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]]
+// CK24-DAG: [[CBPVAL1]] = bitcast [[SA]]** [[SEC0]] to i8*
+// CK24-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8*
+// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SA]]* [[SEC11:%[^,]+]], i{{.+}} 0
+// CK24-DAG: [[SEC11]] = load [[SA]]*, [[SA]]** [[SEC111:%[^,]+]],
+// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[[SB]]* [[SEC1111:[^,]+]], i{{.+}} 0, i{{.+}} 4
+// CK24-DAG: [[SEC1111]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1
+
+// CK24: call void [[CALL09:@.+]]([[SC]]* {{[^,]+}}) 
+#pragma omp target map(s.s.p->a)
+  { s.a++; }
+
+// Region 10
+// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE10]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE10]]{{.+}})
+// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8*
+// CK24-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
+// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[10 x i32]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SA]]* [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 2
+// CK24-DAG: [[SEC000]] = getelementptr {{.*}}[[SB]]* [[SEC0000:%[^,]+]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: [[SEC0000]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1
+
+// CK24: call void [[CALL10:@.+]]([[SC]]* {{[^,]+}})
+#pragma omp target map(s.s.s.b[:2])
+  { s.a++; }
+  
+// Region 11
+// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE11]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE11]]{{.+}})
+// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8*
+// CK24-DAG: [[CPVAL0]] = bitcast [[SA]]** [[SEC0:%.+]] to i8*
+// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SB]]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 4
+// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1
+
+// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]]
+// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]]
+// CK24-DAG: [[CBPVAL1]] = bitcast [[SA]]** [[SEC0]] to i8*
+// CK24-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8*
+// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[10 x i32]* [[SEC11:%[^,]+]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[SEC11]] = getelementptr {{.*}}[[SA]]* [[SEC111:%[^,]+]], i{{.+}} 0, i{{.+}} 2
+// CK24-DAG: [[SEC111]] = load [[SA]]*, [[SA]]** [[SEC1111:%[^,]+]],
+// CK24-DAG: [[SEC1111]] = getelementptr {{.*}}[[SB]]* [[SEC11111:%[^,]+]], i{{.+}} 0, i{{.+}} 4
+// CK24-DAG: [[SEC11111]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1
+
+// CK24: call void [[CALL11:@.+]]([[SC]]* {{[^,]+}}) 
+#pragma omp target map(s.s.p->b[:2])
+  { s.a++; }
+
+// Region 12
+// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 4, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[4 x i{{.+}}]* [[SIZE12]], {{.+}}getelementptr {{.+}}[4 x i{{.+}}]* [[MTYPE12]]{{.+}})
+// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8*
+// CK24-DAG: [[CPVAL0]] = bitcast [[SB]]** [[SEC0:%.+]] to i8*
+// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2
+
+// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]]
+// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]]
+// CK24-DAG: [[CBPVAL1]] = bitcast [[SB]]** [[SEC0]] to i8*
+// CK24-DAG: [[CPVAL1]] = bitcast [[SA]]** [[SEC1:%.+]] to i8*
+// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SB]]* [[SEC11:%[^,]+]], i{{.+}} 0, i{{.+}} 4
+// CK24-DAG: [[SEC11]] = load [[SB]]*, [[SB]]** [[SEC111:%[^,]+]],
+// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2
+
+// CK24-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
+// CK24-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
+// CK24-DAG: store i8* [[CBPVAL2:%[^,]+]], i8** [[BP2]]
+// CK24-DAG: store i8* [[CPVAL2:%[^,]+]], i8** [[P2]]
+// CK24-DAG: [[CBPVAL2]] = bitcast [[SA]]** [[SEC1]] to i8*
+// CK24-DAG: [[CPVAL2]] = bitcast [[SA]]** [[SEC2:%.+]] to i8*
+// CK24-DAG: [[SEC2]] = getelementptr {{.*}}[[SA]]* [[SEC22:%[^,]+]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: [[SEC22]] = load [[SA]]*, [[SA]]** [[SEC222:%[^,]+]],
+// CK24-DAG: [[SEC222]] = getelementptr {{.*}}[[SB]]* [[SEC2222:%[^,]+]], i{{.+}} 0, i{{.+}} 4
+// CK24-DAG: [[SEC2222]] = load [[SB]]*, [[SB]]** [[SEC22222:%[^,]+]],
+// CK24-DAG: [[SEC22222]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2
+
+// CK24-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3
+// CK24-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3
+// CK24-DAG: store i8* [[CBPVAL3:%[^,]+]], i8** [[BP3]]
+// CK24-DAG: store i8* [[CPVAL3:%[^,]+]], i8** [[P3]]
+// CK24-DAG: [[CBPVAL3]] = bitcast [[SA]]** [[SEC2]] to i8*
+// CK24-DAG: [[CPVAL3]] = bitcast i32* [[SEC3:%.+]] to i8*
+// CK24-DAG: [[SEC3]] = getelementptr {{.*}}[[SA]]* [[SEC33:%[^,]+]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[SEC33]] = load [[SA]]*, [[SA]]** [[SEC333:%[^,]+]],
+// CK24-DAG: [[SEC333]] = getelementptr {{.*}}[[SA]]* [[SEC3333:%[^,]+]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: [[SEC3333]] = load [[SA]]*, [[SA]]** [[SEC33333:%[^,]+]],
+// CK24-DAG: [[SEC33333]] = getelementptr {{.*}}[[SB]]* [[SEC333333:%[^,]+]], i{{.+}} 0, i{{.+}} 4
+// CK24-DAG: [[SEC333333]] = load [[SB]]*, [[SB]]** [[SEC3333333:%[^,]+]],
+// CK24-DAG: [[SEC3333333]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2
+
+// CK24: call void [[CALL12:@.+]]([[SC]]* {{[^,]+}})
+#pragma omp target map(s.p->p->p->a)
+  { s.a++; }
+
+//
+// Same thing but starting from a pointer.
+//
+// Region 13
+// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE13]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE13]]{{.+}})
+// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8*
+// CK24-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
+// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 0
+
+// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}}
+// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}}
+
+// CK24: call void [[CALL13:@.+]]([[SC]]* {{[^,]+}})
+#pragma omp target map(p->a)
+  { p->a++; }
+  
+// Region 14
+// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE14]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE14]]{{.+}})
+// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8*
+// CK24-DAG: [[CPVAL0]] = bitcast [[SA]]* [[SEC0:%.+]] to i8*
+// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SB]]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1
+
+// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}}
+// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}}
+
+// CK24: call void [[CALL14:@.+]]([[SC]]* {{[^,]+}})
+#pragma omp target map(p->s.s)
+  { p->a++; }
+  
+// Region 15
+// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE15]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE15]]{{.+}})
+// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8*
+// CK24-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
+// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SA]]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SB]]* [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: [[SEC000]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1
+
+// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}}
+// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}}
+
+// CK24: call void [[CALL15:@.+]]([[SC]]* {{[^,]+}})
+#pragma omp target map(p->s.s.a)
+  { p->a++; }
+
+// Region 16
+// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE16]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE16]]{{.+}})
+// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8*
+// CK24-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
+// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[10 x i32]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 3
+
+// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}}
+// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}}
+
+// CK24: call void [[CALL16:@.+]]([[SC]]* {{[^,]+}})
+#pragma omp target map(p->b[:5])
+  { p->a++; }
+  
+// Region 17
+// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE17]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE17]]{{.+}})
+// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8*
+// CK24-DAG: [[CPVAL0]] = bitcast [[SB]]** [[SEC0:%.+]] to i8*
+// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 2
+
+// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]]
+// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]]
+// CK24-DAG: [[CBPVAL1]] = bitcast [[SB]]** [[SEC0]] to i8*
+// CK24-DAG: [[CPVAL1]] = bitcast [[SB]]* [[SEC1:%.+]] to i8*
+// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SB]]* [[SEC11:%[^,]+]], i{{.+}} 0
+// CK24-DAG: [[SEC11]] = load [[SB]]*, [[SB]]** [[SEC111:%[^,]+]],
+// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[[SC]]* [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 2
+
+// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}}
+// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}}
+// CK24-DAG: [[VAR000]] = load [[SC]]*, [[SC]]** %{{.+}}
+
+// CK24: call void [[CALL17:@.+]]([[SC]]* {{[^,]+}}) 
+#pragma omp target map(p->p[:5])
+  { p->a++; }
+
+// Region 18
+// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE18]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE18]]{{.+}})
+// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8*
+// CK24-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
+// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SA]]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[10 x [[SA]]]* [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 3
+// CK24-DAG: [[SEC000]] = getelementptr {{.*}}[[SB]]* [[SEC0000:%[^,]+]], i{{.+}} 0, i{{.+}} 2
+// CK24-DAG: [[SEC0000]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1
+
+// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}}
+// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}}
+
+// CK24: call void [[CALL18:@.+]]([[SC]]* {{[^,]+}})
+#pragma omp target map(p->s.sa[3].a)
+  { p->a++; }
+  
+// Region 19
+// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE19]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE19]]{{.+}})
+// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8*
+// CK24-DAG: [[CPVAL0]] = bitcast [[SA]]** [[SEC0:%.+]] to i8*
+// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[10 x [[SA]]*]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 3
+// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SB]]* [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 3
+// CK24-DAG: [[SEC000]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1
+
+// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]]
+// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]]
+// CK24-DAG: [[CBPVAL1]] = bitcast [[SA]]** [[SEC0]] to i8*
+// CK24-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8*
+// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SA]]* [[SEC11:%[^,]+]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[SEC11]] = load [[SA]]*, [[SA]]** [[SEC111:%[^,]+]],
+// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[10 x [[SA]]*]* [[SEC1111:%[^,]+]], i{{.+}} 0, i{{.+}} 3
+// CK24-DAG: [[SEC1111]] = getelementptr {{.*}}[[SB]]* [[SEC11111:%[^,]+]], i{{.+}} 0, i{{.+}} 3
+// CK24-DAG: [[SEC11111]] = getelementptr {{.*}}[[SC]]* [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 1
+
+// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}}
+// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}}
+// CK24-DAG: [[VAR000]] = load [[SC]]*, [[SC]]** %{{.+}}
+
+// CK24: call void [[CALL19:@.+]]([[SC]]* {{[^,]+}}) 
+#pragma omp target map(p->s.sp[3]->a)
+  { p->a++; }
+
+// Region 20
+// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE20]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE20]]{{.+}})
+// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8*
+// CK24-DAG: [[CPVAL0]] = bitcast [[SB]]** [[SEC0:%.+]] to i8*
+// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 2
+
+// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]]
+// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]]
+// CK24-DAG: [[CBPVAL1]] = bitcast [[SB]]** [[SEC0]] to i8*
+// CK24-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8*
+// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SB]]* [[SEC11:%[^,]+]], i{{.+}} 0
+// CK24-DAG: [[SEC11]] = load [[SB]]*, [[SB]]** [[SEC111:%[^,]+]],
+// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[[SC]]* [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 2
+
+// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}}
+// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}}
+// CK24-DAG: [[VAR000]] = load [[SC]]*, [[SC]]** %{{.+}}
+
+// CK24: call void [[CALL20:@.+]]([[SC]]* {{[^,]+}})
+#pragma omp target map(p->p->a)
+  { p->a++; }
+  
+// Region 21
+// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE21]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE21]]{{.+}})
+// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8*
+// CK24-DAG: [[CPVAL0]] = bitcast [[SA]]** [[SEC0:%.+]] to i8*
+// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SB]]* [[SEC00:[^,]+]], i{{.+}} 0, i{{.+}} 4
+// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1
+
+// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]]
+// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]]
+// CK24-DAG: [[CBPVAL1]] = bitcast [[SA]]** [[SEC0]] to i8*
+// CK24-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8*
+// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SA]]* [[SEC11:%[^,]+]], i{{.+}} 0
+// CK24-DAG: [[SEC11]] = load [[SA]]*, [[SA]]** [[SEC111:%[^,]+]],
+// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[[SB]]* [[SEC1111:[^,]+]], i{{.+}} 0, i{{.+}} 4
+// CK24-DAG: [[SEC1111]] = getelementptr {{.*}}[[SC]]* [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 1
+
+// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}}
+// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}}
+// CK24-DAG: [[VAR000]] = load [[SC]]*, [[SC]]** %{{.+}}
+
+// CK24: call void [[CALL21:@.+]]([[SC]]* {{[^,]+}}) 
+#pragma omp target map(p->s.p->a)
+  { p->a++; }
+
+// Region 22
+// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE22]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE22]]{{.+}})
+// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8*
+// CK24-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8*
+// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[10 x i32]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SA]]* [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 2
+// CK24-DAG: [[SEC000]] = getelementptr {{.*}}[[SB]]* [[SEC0000:%[^,]+]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: [[SEC0000]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1
+
+// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}}
+// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}}
+
+// CK24: call void [[CALL22:@.+]]([[SC]]* {{[^,]+}})
+#pragma omp target map(p->s.s.b[:2])
+  { p->a++; }
+  
+// Region 23
+// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE23]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE23]]{{.+}})
+// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8*
+// CK24-DAG: [[CPVAL0]] = bitcast [[SA]]** [[SEC0:%.+]] to i8*
+// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SB]]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 4
+// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1
+
+// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]]
+// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]]
+// CK24-DAG: [[CBPVAL1]] = bitcast [[SA]]** [[SEC0]] to i8*
+// CK24-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8*
+// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[10 x i32]* [[SEC11:%[^,]+]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[SEC11]] = getelementptr {{.*}}[[SA]]* [[SEC111:%[^,]+]], i{{.+}} 0, i{{.+}} 2
+// CK24-DAG: [[SEC111]] = load [[SA]]*, [[SA]]** [[SEC1111:%[^,]+]],
+// CK24-DAG: [[SEC1111]] = getelementptr {{.*}}[[SB]]* [[SEC11111:%[^,]+]], i{{.+}} 0, i{{.+}} 4
+// CK24-DAG: [[SEC11111]] = getelementptr {{.*}}[[SC]]* [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 1
+
+// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}}
+// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}}
+// CK24-DAG: [[VAR000]] = load [[SC]]*, [[SC]]** %{{.+}}
+
+// CK24: call void [[CALL23:@.+]]([[SC]]* {{[^,]+}}) 
+#pragma omp target map(p->s.p->b[:2])
+  { p->a++; }
+
+// Region 24
+// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 4, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[4 x i{{.+}}]* [[SIZE24]], {{.+}}getelementptr {{.+}}[4 x i{{.+}}]* [[MTYPE24]]{{.+}})
+// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8*
+// CK24-DAG: [[CPVAL0]] = bitcast [[SB]]** [[SEC0:%.+]] to i8*
+// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 2
+
+// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]]
+// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]]
+// CK24-DAG: [[CBPVAL1]] = bitcast [[SB]]** [[SEC0]] to i8*
+// CK24-DAG: [[CPVAL1]] = bitcast [[SA]]** [[SEC1:%.+]] to i8*
+// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SB]]* [[SEC11:%[^,]+]], i{{.+}} 0, i{{.+}} 4
+// CK24-DAG: [[SEC11]] = load [[SB]]*, [[SB]]** [[SEC111:%[^,]+]],
+// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[[SC]]* [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 2
+
+// CK24-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
+// CK24-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
+// CK24-DAG: store i8* [[CBPVAL2:%[^,]+]], i8** [[BP2]]
+// CK24-DAG: store i8* [[CPVAL2:%[^,]+]], i8** [[P2]]
+// CK24-DAG: [[CBPVAL2]] = bitcast [[SA]]** [[SEC1]] to i8*
+// CK24-DAG: [[CPVAL2]] = bitcast [[SA]]** [[SEC2:%.+]] to i8*
+// CK24-DAG: [[SEC2]] = getelementptr {{.*}}[[SA]]* [[SEC22:%[^,]+]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: [[SEC22]] = load [[SA]]*, [[SA]]** [[SEC222:%[^,]+]],
+// CK24-DAG: [[SEC222]] = getelementptr {{.*}}[[SB]]* [[SEC2222:%[^,]+]], i{{.+}} 0, i{{.+}} 4
+// CK24-DAG: [[SEC2222]] = load [[SB]]*, [[SB]]** [[SEC22222:%[^,]+]],
+// CK24-DAG: [[SEC22222]] = getelementptr {{.*}}[[SC]]* [[VAR0000:%.+]], i{{.+}} 0, i{{.+}} 2
+
+// CK24-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3
+// CK24-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3
+// CK24-DAG: store i8* [[CBPVAL3:%[^,]+]], i8** [[BP3]]
+// CK24-DAG: store i8* [[CPVAL3:%[^,]+]], i8** [[P3]]
+// CK24-DAG: [[CBPVAL3]] = bitcast [[SA]]** [[SEC2]] to i8*
+// CK24-DAG: [[CPVAL3]] = bitcast i32* [[SEC3:%.+]] to i8*
+// CK24-DAG: [[SEC3]] = getelementptr {{.*}}[[SA]]* [[SEC33:%[^,]+]], i{{.+}} 0, i{{.+}} 0
+// CK24-DAG: [[SEC33]] = load [[SA]]*, [[SA]]** [[SEC333:%[^,]+]],
+// CK24-DAG: [[SEC333]] = getelementptr {{.*}}[[SA]]* [[SEC3333:%[^,]+]], i{{.+}} 0, i{{.+}} 1
+// CK24-DAG: [[SEC3333]] = load [[SA]]*, [[SA]]** [[SEC33333:%[^,]+]],
+// CK24-DAG: [[SEC33333]] = getelementptr {{.*}}[[SB]]* [[SEC333333:%[^,]+]], i{{.+}} 0, i{{.+}} 4
+// CK24-DAG: [[SEC333333]] = load [[SB]]*, [[SB]]** [[SEC3333333:%[^,]+]],
+// CK24-DAG: [[SEC3333333]] = getelementptr {{.*}}[[SC]]* [[VAR00000:%.+]], i{{.+}} 0, i{{.+}} 2
+
+// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}}
+// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}}
+// CK24-DAG: [[VAR000]] = load [[SC]]*, [[SC]]** %{{.+}}
+// CK24-DAG: [[VAR0000]] = load [[SC]]*, [[SC]]** %{{.+}}
+// CK24-DAG: [[VAR00000]] = load [[SC]]*, [[SC]]** %{{.+}}
+
+// CK24: call void [[CALL24:@.+]]([[SC]]* {{[^,]+}})
+#pragma omp target map(p->p->p->p->a)
+  { p->a++; }
+
+  return s.a;
+}
+
+// CK24: define {{.+}}[[CALL01]]
+// CK24: define {{.+}}[[CALL02]]
+// CK24: define {{.+}}[[CALL03]]
+// CK24: define {{.+}}[[CALL04]]
+// CK24: define {{.+}}[[CALL05]]
+// CK24: define {{.+}}[[CALL06]]
+// CK24: define {{.+}}[[CALL07]]
+// CK24: define {{.+}}[[CALL08]]
+// CK24: define {{.+}}[[CALL09]]
+// CK24: define {{.+}}[[CALL10]]
+// CK24: define {{.+}}[[CALL11]]
+// CK24: define {{.+}}[[CALL12]]
+// CK24: define {{.+}}[[CALL13]]
+// CK24: define {{.+}}[[CALL14]]
+// CK24: define {{.+}}[[CALL15]]
+// CK24: define {{.+}}[[CALL16]]
+// CK24: define {{.+}}[[CALL17]]
+// CK24: define {{.+}}[[CALL18]]
+// CK24: define {{.+}}[[CALL19]]
+// CK24: define {{.+}}[[CALL20]]
+// CK24: define {{.+}}[[CALL21]]
+// CK24: define {{.+}}[[CALL22]]
+// CK24: define {{.+}}[[CALL23]]
+// CK24: define {{.+}}[[CALL24]]
+#endif
+
 #endif