Page MenuHomePhabricator

[OpenMP] Map clause codegeneration.
ClosedPublic

Authored by sfantao on Jan 29 2016, 11:58 PM.

Details

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.

Diff Detail

Event Timeline

sfantao updated this revision to Diff 46469.Jan 29 2016, 11:58 PM
sfantao retitled this revision from to [OpenMP] Map clause codegeneration..
sfantao updated this object.
ABataev added inline comments.Feb 15 2016, 12:01 AM
lib/CodeGen/CGExpr.cpp
2937–2943 ↗(On Diff #46469)

This must be fixed already, please reabse

lib/CodeGen/CGOpenMPRuntime.cpp
4752–4783

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'.

4963–4965

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 updated this revision to Diff 48085.Feb 16 2016, 10:54 AM
sfantao marked 2 inline comments as done.

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

Hi Alexey,

Thanks for the review!

lib/CodeGen/CGExpr.cpp
2937–2943 ↗(On Diff #46469)

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

lib/CodeGen/CGOpenMPRuntime.cpp
4752–4783

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.

4963–4965

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

ABataev added inline comments.Feb 16 2016, 8:34 PM
lib/CodeGen/CGOpenMPRuntime.cpp
4752–4783

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

sfantao marked an inline comment as done.Feb 17 2016, 7:55 AM

Hi Alexey,

lib/CodeGen/CGOpenMPRuntime.cpp
4752–4783

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

sfantao marked an inline comment as done.Feb 17 2016, 7:57 AM
sfantao added inline comments.
test/OpenMP/target_map_codegen.cpp
1705

Regression tests for array sections start here.

ABataev added inline comments.Feb 18 2016, 4:28 AM
lib/CodeGen/CGOpenMPRuntime.cpp
4752–4783

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 updated this revision to Diff 48365.Feb 18 2016, 11:02 AM

Add regression tests for invalid array sections and multidimensional VLA.

Hi Alexey,

Thanks for the feedback!

lib/CodeGen/CGOpenMPRuntime.cpp
4752–4783

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 added inline comments.Feb 18 2016, 11:20 AM
lib/CodeGen/CGOpenMPRuntime.cpp
4752–4783

'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.

Hi Alexey,

lib/CodeGen/CGOpenMPRuntime.cpp
4752–4783

'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 added inline comments.Feb 19 2016, 4:43 AM
lib/CodeGen/CGOpenMPRuntime.cpp
4752–4783

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 updated this revision to Diff 48830.Feb 23 2016, 10:19 AM

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

lib/CodeGen/CGOpenMPRuntime.cpp
4752–4783

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 updated this revision to Diff 49240.Feb 26 2016, 2:25 PM

Rebase and use getBaseOriginalType.

sfantao updated this revision to Diff 49711.Mar 2 2016, 10:22 PM
  • Rebase.
  • Add regression test for non-contiguous storage map.
sfantao updated this revision to Diff 50148.Mar 9 2016, 8:57 AM

Rebase.

ABataev edited edge metadata.Mar 17 2016, 5:23 AM

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) {

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

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

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

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

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 updated this revision to Diff 51204.Mar 21 2016, 11:48 AM
sfantao edited edge metadata.
  • Add regression tests for inner captures in the presence of map clause.
  • Rebase.
ABataev added inline comments.Apr 4 2016, 1:54 AM
lib/CodeGen/CGOpenMPRuntime.cpp
4690–4728

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?

4759

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

4790–4795

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

4978–4981

CGF has EmitLoadOfPointer() function

4996–5032

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

4997

auto &&

5091

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

sfantao updated this revision to Diff 54558.Apr 21 2016, 12:00 PM
sfantao marked 7 inline comments as done.
  • 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.

Hi Alexey,

Thanks for the review!

lib/CodeGen/CGOpenMPRuntime.cpp
4690–4728

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.

4759

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

4790–4795

Ok, I removed this.

4978–4981

Ok, using EmitLoadOfPointerLValue now.

4996–5032

Moved this code to a new member function as suggested.

4997

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

5091

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 updated this revision to Diff 54673.Apr 22 2016, 9:54 AM
  • Rebase.
sfantao updated this revision to Diff 54942.Apr 25 2016, 4:27 PM
  • Use canonical declaration when dealng with mappable expressions.
sfantao updated this revision to Diff 55009.Apr 26 2016, 8:29 AM
  • Rebase.
ABataev accepted this revision.Apr 27 2016, 1:09 AM
ABataev edited edge metadata.
ABataev added inline comments.
lib/CodeGen/CGOpenMPRuntime.cpp
4699–4701

Remove braces

5093

cast is not required, remove it

This revision is now accepted and ready to land.Apr 27 2016, 1:09 AM
sfantao updated this revision to Diff 55343.Apr 27 2016, 3:44 PM
sfantao marked 2 inline comments as done.
sfantao edited edge metadata.
  • Remove extra braces.

Hi Alexey,

Thanks for the review!

lib/CodeGen/CGOpenMPRuntime.cpp
4699–4701

Done!

5093

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.

sfantao closed this revision.Apr 27 2016, 3:46 PM