[OpenMP] Parsing and sema support for target update directive

Summary:
This patch is to add parsing and sema support for `target update` directive. Support for the `to` and `from` clauses will be added by a different patch.  This patch also adds support for other clauses that are already implemented upstream and apply to `target update`, e.g. `device` and `if`.

This patch is based on the original post by Kelvin Li.

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

Subscribers: caomhin, cfe-commits

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

llvm-svn: 270878

clang/include/clang-c/Index.h

@@ -2281,7 +2281,7 @@ enum CXCursorKind {
    */
   CXCursor_OMPTaskLoopSimdDirective      = 259,
 
-   /** \brief OpenMP distribute directive.
+  /** \brief OpenMP distribute directive.
    */
   CXCursor_OMPDistributeDirective        = 260,
 
@@ -2301,7 +2301,11 @@ enum CXCursorKind {
    */
   CXCursor_OMPTargetParallelForDirective = 264,
 
-  CXCursor_LastStmt                   = CXCursor_OMPTargetParallelForDirective,
+  /** \brief OpenMP target update directive.
+   */
+  CXCursor_OMPTargetUpdateDirective      = 265,
+
+  CXCursor_LastStmt                      = CXCursor_OMPTargetUpdateDirective,
 
   /**
    * \brief Cursor that represents the translation unit itself.

clang/include/clang/AST/RecursiveASTVisitor.h

@@ -2502,6 +2502,9 @@ DEF_TRAVERSE_STMT(OMPTargetParallelForDirective,
 DEF_TRAVERSE_STMT(OMPTeamsDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
+DEF_TRAVERSE_STMT(OMPTargetUpdateDirective,
+                  { TRY_TO(TraverseOMPExecutableDirective(S)); })
+
 DEF_TRAVERSE_STMT(OMPTaskLoopDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 

clang/include/clang/AST/StmtOpenMP.h

@@ -2688,7 +2688,7 @@ class OMPDistributeDirective : public OMPLoopDirective {
   /// \param Clauses List of clauses.
   /// \param AssociatedStmt Statement, associated with the directive.
   /// \param Exprs Helper expressions for CodeGen.
-    ///
+  ///
   static OMPDistributeDirective *
   Create(const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
          unsigned CollapsedNum, ArrayRef<OMPClause *> Clauses,
@@ -2710,6 +2710,65 @@ class OMPDistributeDirective : public OMPLoopDirective {
   }
 };
 
+/// \brief This represents '#pragma omp target update' directive.
+///
+/// \code
+/// #pragma omp target update to(a) from(b) device(1)
+/// \endcode
+/// In this example directive '#pragma omp target update' has clause 'to' with
+/// argument 'a', clause 'from' with argument 'b' and clause 'device' with
+/// argument '1'.
+///
+class OMPTargetUpdateDirective : public OMPExecutableDirective {
+  friend class ASTStmtReader;
+  /// \brief Build directive with the given start and end location.
+  ///
+  /// \param StartLoc Starting location of the directive kind.
+  /// \param EndLoc Ending Location of the directive.
+  /// \param NumClauses The number of clauses.
+  ///
+  OMPTargetUpdateDirective(SourceLocation StartLoc, SourceLocation EndLoc,
+                           unsigned NumClauses)
+      : OMPExecutableDirective(this, OMPTargetUpdateDirectiveClass,
+                               OMPD_target_update, StartLoc, EndLoc, NumClauses,
+                               0) {}
+
+  /// \brief Build an empty directive.
+  ///
+  /// \param NumClauses Number of clauses.
+  ///
+  explicit OMPTargetUpdateDirective(unsigned NumClauses)
+      : OMPExecutableDirective(this, OMPTargetUpdateDirectiveClass,
+                               OMPD_target_update, SourceLocation(),
+                               SourceLocation(), NumClauses, 0) {}
+
+public:
+  /// \brief Creates directive with a list of \a Clauses.
+  ///
+  /// \param C AST context.
+  /// \param StartLoc Starting location of the directive kind.
+  /// \param EndLoc Ending Location of the directive.
+  /// \param Clauses List of clauses.
+  ///
+  static OMPTargetUpdateDirective *Create(const ASTContext &C,
+                                          SourceLocation StartLoc,
+                                          SourceLocation EndLoc,
+                                          ArrayRef<OMPClause *> Clauses);
+
+  /// \brief Creates an empty directive with the place for \a NumClauses
+  /// clauses.
+  ///
+  /// \param C AST context.
+  /// \param NumClauses The number of clauses.
+  ///
+  static OMPTargetUpdateDirective *CreateEmpty(const ASTContext &C,
+                                               unsigned NumClauses, EmptyShell);
+
+  static bool classof(const Stmt *T) {
+    return T->getStmtClass() == OMPTargetUpdateDirectiveClass;
+  }
+};
+
 } // end namespace clang
 
 #endif

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -8265,6 +8265,8 @@ def err_omp_expected_uniform_param : Error<
   "expected a reference to a parameter specified in a 'uniform' clause">;
 def err_omp_expected_int_param : Error<
   "expected a reference to an integer-typed parameter">;
+def err_omp_at_least_one_motion_clause_required : Error<
+  "expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'">;
 } // end of OpenMP category
 
 let CategoryName = "Related Result Type Issue" in {

clang/include/clang/Basic/OpenMPKinds.def

@@ -72,6 +72,9 @@
 #ifndef OPENMP_TARGET_PARALLEL_FOR_CLAUSE
 #  define OPENMP_TARGET_PARALLEL_FOR_CLAUSE(Name)
 #endif
+#ifndef OPENMP_TARGET_UPDATE_CLAUSE
+#  define OPENMP_TARGET_UPDATE_CLAUSE(Name)
+#endif
 #ifndef OPENMP_TEAMS_CLAUSE
 #  define OPENMP_TEAMS_CLAUSE(Name)
 #endif
@@ -150,6 +153,7 @@ OPENMP_DIRECTIVE_EXT(target_enter_data, "target enter data")
 OPENMP_DIRECTIVE_EXT(target_exit_data, "target exit data")
 OPENMP_DIRECTIVE_EXT(target_parallel, "target parallel")
 OPENMP_DIRECTIVE_EXT(target_parallel_for, "target parallel for")
+OPENMP_DIRECTIVE_EXT(target_update, "target update")
 OPENMP_DIRECTIVE_EXT(parallel_for, "parallel for")
 OPENMP_DIRECTIVE_EXT(parallel_for_simd, "parallel for simd")
 OPENMP_DIRECTIVE_EXT(parallel_sections, "parallel sections")
@@ -442,6 +446,11 @@ OPENMP_TARGET_PARALLEL_FOR_CLAUSE(schedule)
 OPENMP_TARGET_PARALLEL_FOR_CLAUSE(ordered)
 OPENMP_TARGET_PARALLEL_FOR_CLAUSE(linear)
 
+// Clauses allowed for OpenMP directive 'target update'.
+// TODO More clauses for 'target update' directive.
+OPENMP_TARGET_UPDATE_CLAUSE(if)
+OPENMP_TARGET_UPDATE_CLAUSE(device)
+
 // Clauses allowed for OpenMP directive 'teams'.
 // TODO More clauses for 'teams' directive.
 OPENMP_TEAMS_CLAUSE(default)
@@ -553,3 +562,4 @@ OPENMP_DIST_SCHEDULE_KIND(static)
 #undef OPENMP_DIST_SCHEDULE_KIND
 #undef OPENMP_DEFAULTMAP_KIND
 #undef OPENMP_DEFAULTMAP_MODIFIER
+#undef OPENMP_TARGET_UPDATE_CLAUSE

clang/include/clang/Basic/StmtNodes.td

@@ -220,6 +220,7 @@ def OMPTargetEnterDataDirective : DStmt<OMPExecutableDirective>;
 def OMPTargetExitDataDirective : DStmt<OMPExecutableDirective>;
 def OMPTargetParallelDirective : DStmt<OMPExecutableDirective>;
 def OMPTargetParallelForDirective : DStmt<OMPExecutableDirective>;
+def OMPTargetUpdateDirective : DStmt<OMPExecutableDirective>;
 def OMPTeamsDirective : DStmt<OMPExecutableDirective>;
 def OMPCancellationPointDirective : DStmt<OMPExecutableDirective>;
 def OMPCancelDirective : DStmt<OMPExecutableDirective>;

clang/include/clang/Sema/Sema.h

@@ -8173,6 +8173,10 @@ class Sema {
       ArrayRef<OMPClause *> Clauses, Stmt *AStmt, SourceLocation StartLoc,
       SourceLocation EndLoc,
       llvm::DenseMap<ValueDecl *, Expr *> &VarsWithImplicitDSA);
+  /// \brief Called on well-formed '\#pragma omp target update'.
+  StmtResult ActOnOpenMPTargetUpdateDirective(ArrayRef<OMPClause *> Clauses,
+                                              SourceLocation StartLoc,
+                                              SourceLocation EndLoc);
 
   /// Checks correctness of linear modifiers.
   bool CheckOpenMPLinearModifier(OpenMPLinearClauseKind LinKind,

clang/include/clang/Serialization/ASTBitCodes.h

@@ -1465,6 +1465,7 @@ namespace clang {
       STMT_OMP_TASKLOOP_DIRECTIVE,
       STMT_OMP_TASKLOOP_SIMD_DIRECTIVE,
       STMT_OMP_DISTRIBUTE_DIRECTIVE,
+      STMT_OMP_TARGET_UPDATE_DIRECTIVE,
       EXPR_OMP_ARRAY_SECTION,
 
       // ARC

clang/lib/AST/StmtOpenMP.cpp

@@ -1012,3 +1012,25 @@ OMPDistributeDirective::CreateEmpty(const ASTContext &C, unsigned NumClauses,
                              numLoopChildren(CollapsedNum, OMPD_distribute));
   return new (Mem) OMPDistributeDirective(CollapsedNum, NumClauses);
 }
+
+OMPTargetUpdateDirective *
+OMPTargetUpdateDirective::Create(const ASTContext &C, SourceLocation StartLoc,
+                                 SourceLocation EndLoc,
+                                 ArrayRef<OMPClause *> Clauses) {
+  unsigned Size = llvm::alignTo(sizeof(OMPTargetUpdateDirective),
+                                llvm::alignOf<OMPClause *>());
+  void *Mem = C.Allocate(Size + sizeof(OMPClause *) * Clauses.size());
+  OMPTargetUpdateDirective *Dir =
+      new (Mem) OMPTargetUpdateDirective(StartLoc, EndLoc, Clauses.size());
+  Dir->setClauses(Clauses);
+  return Dir;
+}
+
+OMPTargetUpdateDirective *
+OMPTargetUpdateDirective::CreateEmpty(const ASTContext &C, unsigned NumClauses,
+                                      EmptyShell) {
+  unsigned Size = llvm::alignTo(sizeof(OMPTargetUpdateDirective),
+                                llvm::alignOf<OMPClause *>());
+  void *Mem = C.Allocate(Size + sizeof(OMPClause *) * NumClauses);
+  return new (Mem) OMPTargetUpdateDirective(NumClauses);
+}

clang/lib/AST/StmtPrinter.cpp

@@ -1131,6 +1131,12 @@ void StmtPrinter::VisitOMPDistributeDirective(OMPDistributeDirective *Node) {
   PrintOMPExecutableDirective(Node);
 }
 
+void StmtPrinter::VisitOMPTargetUpdateDirective(
+    OMPTargetUpdateDirective *Node) {
+  Indent() << "#pragma omp target update ";
+  PrintOMPExecutableDirective(Node);
+}
+
 //===----------------------------------------------------------------------===//
 //  Expr printing methods.
 //===----------------------------------------------------------------------===//

clang/lib/AST/StmtProfile.cpp

@@ -656,6 +656,11 @@ void OMPClauseProfiler::VisitOMPDistScheduleClause(
 
 void OMPClauseProfiler::VisitOMPDefaultmapClause(const OMPDefaultmapClause *) {}
 
+void StmtProfiler::VisitOMPTargetUpdateDirective(
+    const OMPTargetUpdateDirective *S) {
+  VisitOMPExecutableDirective(S);
+}
+
 void StmtProfiler::VisitExpr(const Expr *S) {
   VisitStmt(S);
 }

clang/lib/Basic/OpenMPKinds.cpp

@@ -475,6 +475,16 @@ bool clang::isAllowedClauseForDirective(OpenMPDirectiveKind DKind,
 #define OPENMP_TARGET_PARALLEL_FOR_CLAUSE(Name)                                \
   case OMPC_##Name:                                                            \
     return true;
+#include "clang/Basic/OpenMPKinds.def"
+    default:
+      break;
+    }
+    break;
+  case OMPD_target_update:
+    switch (CKind) {
+#define OPENMP_TARGET_UPDATE_CLAUSE(Name)                                      \
+  case OMPC_##Name:                                                            \
+    return true;
 #include "clang/Basic/OpenMPKinds.def"
     default:
       break;
@@ -605,7 +615,7 @@ bool clang::isOpenMPTargetExecutionDirective(OpenMPDirectiveKind DKind) {
 bool clang::isOpenMPTargetDataManagementDirective(OpenMPDirectiveKind DKind) {
   // TODO add target update directive check.
   return DKind == OMPD_target_data || DKind == OMPD_target_enter_data ||
-         DKind == OMPD_target_exit_data;
+         DKind == OMPD_target_exit_data || DKind == OMPD_target_update;
 }
 
 bool clang::isOpenMPTeamsDirective(OpenMPDirectiveKind DKind) {

clang/lib/CodeGen/CGStmt.cpp

@@ -277,6 +277,9 @@ void CodeGenFunction::EmitStmt(const Stmt *S) {
   case Stmt::OMPDistributeDirectiveClass:
     EmitOMPDistributeDirective(cast<OMPDistributeDirective>(*S));
     break;
+  case Stmt::OMPTargetUpdateDirectiveClass:
+    EmitOMPTargetUpdateDirective(cast<OMPTargetUpdateDirective>(*S));
+    break;
   }
 }
 

clang/lib/CodeGen/CGStmtOpenMP.cpp

@@ -3544,3 +3544,9 @@ void CodeGenFunction::EmitOMPTaskLoopSimdDirective(
     const OMPTaskLoopSimdDirective &S) {
   EmitOMPTaskLoopBasedDirective(S);
 }
+
+// Generate the instructions for '#pragma omp target update' directive.
+void CodeGenFunction::EmitOMPTargetUpdateDirective(
+    const OMPTargetUpdateDirective &S) {
+  // TODO: codegen for target update
+}

clang/lib/CodeGen/CodeGenFunction.h

@@ -2380,6 +2380,7 @@ class CodeGenFunction : public CodeGenTypeCache {
   void EmitOMPTargetDataDirective(const OMPTargetDataDirective &S);
   void EmitOMPTargetEnterDataDirective(const OMPTargetEnterDataDirective &S);
   void EmitOMPTargetExitDataDirective(const OMPTargetExitDataDirective &S);
+  void EmitOMPTargetUpdateDirective(const OMPTargetUpdateDirective &S);
   void EmitOMPTargetParallelDirective(const OMPTargetParallelDirective &S);
   void
   EmitOMPTargetParallelForDirective(const OMPTargetParallelForDirective &S);

clang/lib/Parse/ParseOpenMP.cpp

@@ -38,7 +38,8 @@ enum OpenMPDirectiveKindEx {
   OMPD_point,
   OMPD_reduction,
   OMPD_target_enter,
-  OMPD_target_exit
+  OMPD_target_exit,
+  OMPD_update,
 };
 
 class ThreadprivateListParserHelper final {
@@ -73,6 +74,7 @@ static unsigned getOpenMPDirectiveKindEx(StringRef S) {
       .Case("exit", OMPD_exit)
       .Case("point", OMPD_point)
       .Case("reduction", OMPD_reduction)
+      .Case("update", OMPD_update)
       .Default(OMPD_unknown);
 }
 
@@ -90,6 +92,7 @@ static OpenMPDirectiveKind ParseOpenMPDirectiveKind(Parser &P) {
     { OMPD_target, OMPD_data, OMPD_target_data },
     { OMPD_target, OMPD_enter, OMPD_target_enter },
     { OMPD_target, OMPD_exit, OMPD_target_exit },
+    { OMPD_target, OMPD_update, OMPD_target_update },
     { OMPD_target_enter, OMPD_data, OMPD_target_enter_data },
     { OMPD_target_exit, OMPD_data, OMPD_target_exit_data },
     { OMPD_for, OMPD_simd, OMPD_for_simd },
@@ -726,6 +729,7 @@ Parser::DeclGroupPtrTy Parser::ParseOpenMPDeclarativeDirectiveWithExtDecl(
   case OMPD_taskloop_simd:
   case OMPD_distribute:
   case OMPD_end_declare_target:
+  case OMPD_target_update:
     Diag(Tok, diag::err_omp_unexpected_directive)
         << getOpenMPDirectiveName(DKind);
     break;
@@ -756,7 +760,8 @@ Parser::DeclGroupPtrTy Parser::ParseOpenMPDeclarativeDirectiveWithExtDecl(
 ///         'for simd' | 'parallel for simd' | 'target' | 'target data' |
 ///         'taskgroup' | 'teams' | 'taskloop' | 'taskloop simd' |
 ///         'distribute' | 'target enter data' | 'target exit data' |
-///         'target parallel' | 'target parallel for' {clause}
+///         'target parallel' | 'target parallel for' |
+///         'target update' {clause}
 ///         annot_pragma_openmp_end
 ///
 StmtResult Parser::ParseOpenMPDeclarativeOrExecutableDirective(
@@ -830,6 +835,7 @@ StmtResult Parser::ParseOpenMPDeclarativeOrExecutableDirective(
   case OMPD_cancel:
   case OMPD_target_enter_data:
   case OMPD_target_exit_data:
+  case OMPD_target_update:
     if (Allowed == ACK_StatementsOpenMPNonStandalone) {
       Diag(Tok, diag::err_omp_immediate_directive)
           << getOpenMPDirectiveName(DKind) << 0;

clang/lib/Sema/SemaOpenMP.cpp

@@ -1825,6 +1825,7 @@ void Sema::ActOnOpenMPRegionStart(OpenMPDirectiveKind DKind, Scope *CurScope) {
   case OMPD_declare_simd:
   case OMPD_declare_target:
   case OMPD_end_declare_target:
+  case OMPD_target_update:
     llvm_unreachable("OpenMP Directive is not allowed");
   case OMPD_unknown:
     llvm_unreachable("Unknown OpenMP directive");
@@ -3274,6 +3275,12 @@ StmtResult Sema::ActOnOpenMPExecutableDirective(
     Res = ActOnOpenMPDistributeDirective(ClausesWithImplicit, AStmt, StartLoc,
                                          EndLoc, VarsWithInheritedDSA);
     break;
+  case OMPD_target_update:
+    assert(!AStmt && "Statement is not allowed for target update");
+    Res =
+        ActOnOpenMPTargetUpdateDirective(ClausesWithImplicit, StartLoc, EndLoc);
+    AllowedNameModifiers.push_back(OMPD_target_update);
+    break;
   case OMPD_declare_target:
   case OMPD_end_declare_target:
   case OMPD_threadprivate:
@@ -6511,6 +6518,20 @@ Sema::ActOnOpenMPTargetExitDataDirective(ArrayRef<OMPClause *> Clauses,
   return OMPTargetExitDataDirective::Create(Context, StartLoc, EndLoc, Clauses);
 }
 
+StmtResult Sema::ActOnOpenMPTargetUpdateDirective(ArrayRef<OMPClause *> Clauses,
+                                                  SourceLocation StartLoc,
+                                                  SourceLocation EndLoc) {
+  // TODO: Set this flag accordingly when we add support for the 'to' and 'from'
+  // clauses.
+  bool seenMotionClause = false;
+
+  if (!seenMotionClause) {
+    Diag(StartLoc, diag::err_omp_at_least_one_motion_clause_required);
+    return StmtError();
+  }
+  return OMPTargetUpdateDirective::Create(Context, StartLoc, EndLoc, Clauses);
+}
+
 StmtResult Sema::ActOnOpenMPTeamsDirective(ArrayRef<OMPClause *> Clauses,
                                            Stmt *AStmt, SourceLocation StartLoc,
                                            SourceLocation EndLoc) {
@@ -10407,7 +10428,7 @@ Sema::ActOnOpenMPMapClause(OpenMPMapClauseKind MapTypeModifier,
 
     // Save the components and declaration to create the clause. For purposes of
     // the clause creation, any component list that has has base 'this' uses
-    // null has
+    // null as base declaration.
     ClauseComponents.resize(ClauseComponents.size() + 1);
     ClauseComponents.back().append(CurComponents.begin(), CurComponents.end());
     ClauseBaseDeclarations.push_back(isa<MemberExpr>(BE) ? nullptr