[OpenMP] Parsing + sema for target parallel for directive.

Summary:
This patch adds parsing + sema for the target parallel for directive along with testcases.

Reviewers: ABataev

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

llvm-svn: 259654

clang/include/clang-c/Index.h

@@ -2286,7 +2286,11 @@ enum CXCursorKind {
    */
   CXCursor_OMPTargetParallelDirective    = 263,
 
-  CXCursor_LastStmt                      = CXCursor_OMPTargetParallelDirective,
+  /** \brief OpenMP target parallel for directive.
+   */
+  CXCursor_OMPTargetParallelForDirective = 264,
+
+  CXCursor_LastStmt                   = CXCursor_OMPTargetParallelForDirective,
 
   /**
    * \brief Cursor that represents the translation unit itself.

clang/include/clang/AST/RecursiveASTVisitor.h

@@ -2447,6 +2447,9 @@ DEF_TRAVERSE_STMT(OMPTargetExitDataDirective,
 DEF_TRAVERSE_STMT(OMPTargetParallelDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
+DEF_TRAVERSE_STMT(OMPTargetParallelForDirective,
+                  { TRY_TO(TraverseOMPExecutableDirective(S)); })
+
 DEF_TRAVERSE_STMT(OMPTeamsDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 

clang/include/clang/AST/StmtOpenMP.h

@@ -691,7 +691,8 @@ class OMPLoopDirective : public OMPExecutableDirective {
            T->getStmtClass() == OMPParallelForSimdDirectiveClass ||
            T->getStmtClass() == OMPTaskLoopDirectiveClass ||
            T->getStmtClass() == OMPTaskLoopSimdDirectiveClass ||
-           T->getStmtClass() == OMPDistributeDirectiveClass;
+           T->getStmtClass() == OMPDistributeDirectiveClass ||
+           T->getStmtClass() == OMPTargetParallelForDirectiveClass;
   }
 };
 
@@ -2215,6 +2216,87 @@ class OMPTargetParallelDirective : public OMPExecutableDirective {
   }
 };
 
+/// \brief This represents '#pragma omp target parallel for' directive.
+///
+/// \code
+/// #pragma omp target parallel for private(a,b) reduction(+:c,d)
+/// \endcode
+/// In this example directive '#pragma omp target parallel for' has clauses
+/// 'private' with the variables 'a' and 'b' and 'reduction' with operator '+'
+/// and variables 'c' and 'd'.
+///
+class OMPTargetParallelForDirective : public OMPLoopDirective {
+  friend class ASTStmtReader;
+
+  /// \brief true if current region has inner cancel directive.
+  bool HasCancel;
+
+  /// \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 CollapsedNum Number of collapsed nested loops.
+  /// \param NumClauses Number of clauses.
+  ///
+  OMPTargetParallelForDirective(SourceLocation StartLoc, SourceLocation EndLoc,
+                                unsigned CollapsedNum, unsigned NumClauses)
+      : OMPLoopDirective(this, OMPTargetParallelForDirectiveClass,
+                         OMPD_target_parallel_for, StartLoc, EndLoc,
+                         CollapsedNum, NumClauses),
+        HasCancel(false) {}
+
+  /// \brief Build an empty directive.
+  ///
+  /// \param CollapsedNum Number of collapsed nested loops.
+  /// \param NumClauses Number of clauses.
+  ///
+  explicit OMPTargetParallelForDirective(unsigned CollapsedNum,
+                                         unsigned NumClauses)
+      : OMPLoopDirective(this, OMPTargetParallelForDirectiveClass,
+                         OMPD_target_parallel_for, SourceLocation(),
+                         SourceLocation(), CollapsedNum, NumClauses),
+        HasCancel(false) {}
+
+  /// \brief Set cancel state.
+  void setHasCancel(bool Has) { HasCancel = Has; }
+
+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 CollapsedNum Number of collapsed loops.
+  /// \param Clauses List of clauses.
+  /// \param AssociatedStmt Statement, associated with the directive.
+  /// \param Exprs Helper expressions for CodeGen.
+  /// \param HasCancel true if current directive has inner cancel directive.
+  ///
+  static OMPTargetParallelForDirective *
+  Create(const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
+         unsigned CollapsedNum, ArrayRef<OMPClause *> Clauses,
+         Stmt *AssociatedStmt, const HelperExprs &Exprs, bool HasCancel);
+
+  /// \brief Creates an empty directive with the place
+  /// for \a NumClauses clauses.
+  ///
+  /// \param C AST context.
+  /// \param CollapsedNum Number of collapsed nested loops.
+  /// \param NumClauses Number of clauses.
+  ///
+  static OMPTargetParallelForDirective *CreateEmpty(const ASTContext &C,
+                                                    unsigned NumClauses,
+                                                    unsigned CollapsedNum,
+                                                    EmptyShell);
+
+  /// \brief Return true if current directive has inner cancel directive.
+  bool hasCancel() const { return HasCancel; }
+
+  static bool classof(const Stmt *T) {
+    return T->getStmtClass() == OMPTargetParallelForDirectiveClass;
+  }
+};
+
 /// \brief This represents '#pragma omp teams' directive.
 ///
 /// \code

clang/include/clang/Basic/OpenMPKinds.def

@@ -69,6 +69,9 @@
 #ifndef OPENMP_TARGET_PARALLEL_CLAUSE
 #  define OPENMP_TARGET_PARALLEL_CLAUSE(Name)
 #endif
+#ifndef OPENMP_TARGET_PARALLEL_FOR_CLAUSE
+#  define OPENMP_TARGET_PARALLEL_FOR_CLAUSE(Name)
+#endif
 #ifndef OPENMP_TEAMS_CLAUSE
 #  define OPENMP_TEAMS_CLAUSE(Name)
 #endif
@@ -146,6 +149,7 @@ OPENMP_DIRECTIVE_EXT(target_data, "target data")
 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(parallel_for, "parallel for")
 OPENMP_DIRECTIVE_EXT(parallel_for_simd, "parallel for simd")
 OPENMP_DIRECTIVE_EXT(parallel_sections, "parallel sections")
@@ -412,6 +416,27 @@ OPENMP_TARGET_PARALLEL_CLAUSE(proc_bind)
 OPENMP_TARGET_PARALLEL_CLAUSE(shared)
 OPENMP_TARGET_PARALLEL_CLAUSE(reduction)
 
+// Clauses allowed for OpenMP directive 'target parallel for'.
+// TODO: add target clauses 'is_device_ptr'
+OPENMP_TARGET_PARALLEL_FOR_CLAUSE(if)
+OPENMP_TARGET_PARALLEL_FOR_CLAUSE(device)
+OPENMP_TARGET_PARALLEL_FOR_CLAUSE(map)
+OPENMP_TARGET_PARALLEL_FOR_CLAUSE(private)
+OPENMP_TARGET_PARALLEL_FOR_CLAUSE(firstprivate)
+OPENMP_TARGET_PARALLEL_FOR_CLAUSE(lastprivate)
+OPENMP_TARGET_PARALLEL_FOR_CLAUSE(nowait)
+OPENMP_TARGET_PARALLEL_FOR_CLAUSE(depend)
+OPENMP_TARGET_PARALLEL_FOR_CLAUSE(defaultmap)
+OPENMP_TARGET_PARALLEL_FOR_CLAUSE(num_threads)
+OPENMP_TARGET_PARALLEL_FOR_CLAUSE(default)
+OPENMP_TARGET_PARALLEL_FOR_CLAUSE(proc_bind)
+OPENMP_TARGET_PARALLEL_FOR_CLAUSE(shared)
+OPENMP_TARGET_PARALLEL_FOR_CLAUSE(reduction)
+OPENMP_TARGET_PARALLEL_FOR_CLAUSE(collapse)
+OPENMP_TARGET_PARALLEL_FOR_CLAUSE(schedule)
+OPENMP_TARGET_PARALLEL_FOR_CLAUSE(ordered)
+OPENMP_TARGET_PARALLEL_FOR_CLAUSE(linear)
+
 // Clauses allowed for OpenMP directive 'teams'.
 // TODO More clauses for 'teams' directive.
 OPENMP_TEAMS_CLAUSE(default)
@@ -513,6 +538,7 @@ OPENMP_DIST_SCHEDULE_KIND(static)
 #undef OPENMP_TARGET_ENTER_DATA_CLAUSE
 #undef OPENMP_TARGET_EXIT_DATA_CLAUSE
 #undef OPENMP_TARGET_PARALLEL_CLAUSE
+#undef OPENMP_TARGET_PARALLEL_FOR_CLAUSE
 #undef OPENMP_TEAMS_CLAUSE
 #undef OPENMP_SIMD_CLAUSE
 #undef OPENMP_FOR_CLAUSE

clang/include/clang/Basic/StmtNodes.td

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

clang/include/clang/Sema/Sema.h

@@ -7971,6 +7971,12 @@ class Sema {
                                                 Stmt *AStmt,
                                                 SourceLocation StartLoc,
                                                 SourceLocation EndLoc);
+  /// \brief Called on well-formed '\#pragma omp target parallel for' after
+  /// parsing of the  associated statement.
+  StmtResult ActOnOpenMPTargetParallelForDirective(
+      ArrayRef<OMPClause *> Clauses, Stmt *AStmt, SourceLocation StartLoc,
+      SourceLocation EndLoc,
+      llvm::DenseMap<ValueDecl *, Expr *> &VarsWithImplicitDSA);
   /// \brief Called on well-formed '\#pragma omp teams' after parsing of the
   /// associated statement.
   StmtResult ActOnOpenMPTeamsDirective(ArrayRef<OMPClause *> Clauses,

clang/include/clang/Serialization/ASTBitCodes.h

@@ -1451,6 +1451,7 @@ namespace clang {
       STMT_OMP_TARGET_ENTER_DATA_DIRECTIVE,
       STMT_OMP_TARGET_EXIT_DATA_DIRECTIVE,
       STMT_OMP_TARGET_PARALLEL_DIRECTIVE,
+      STMT_OMP_TARGET_PARALLEL_FOR_DIRECTIVE,
       STMT_OMP_TEAMS_DIRECTIVE,
       STMT_OMP_TASKGROUP_DIRECTIVE,
       STMT_OMP_CANCELLATION_POINT_DIRECTIVE,

clang/lib/AST/StmtOpenMP.cpp

@@ -718,6 +718,54 @@ OMPTargetParallelDirective::CreateEmpty(const ASTContext &C,
   return new (Mem) OMPTargetParallelDirective(NumClauses);
 }
 
+OMPTargetParallelForDirective *OMPTargetParallelForDirective::Create(
+    const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
+    unsigned CollapsedNum, ArrayRef<OMPClause *> Clauses, Stmt *AssociatedStmt,
+    const HelperExprs &Exprs, bool HasCancel) {
+  unsigned Size = llvm::alignTo(sizeof(OMPTargetParallelForDirective),
+                                llvm::alignOf<OMPClause *>());
+  void *Mem = C.Allocate(
+      Size + sizeof(OMPClause *) * Clauses.size() +
+      sizeof(Stmt *) * numLoopChildren(CollapsedNum, OMPD_target_parallel_for));
+  OMPTargetParallelForDirective *Dir = new (Mem) OMPTargetParallelForDirective(
+      StartLoc, EndLoc, CollapsedNum, Clauses.size());
+  Dir->setClauses(Clauses);
+  Dir->setAssociatedStmt(AssociatedStmt);
+  Dir->setIterationVariable(Exprs.IterationVarRef);
+  Dir->setLastIteration(Exprs.LastIteration);
+  Dir->setCalcLastIteration(Exprs.CalcLastIteration);
+  Dir->setPreCond(Exprs.PreCond);
+  Dir->setCond(Exprs.Cond);
+  Dir->setInit(Exprs.Init);
+  Dir->setInc(Exprs.Inc);
+  Dir->setIsLastIterVariable(Exprs.IL);
+  Dir->setLowerBoundVariable(Exprs.LB);
+  Dir->setUpperBoundVariable(Exprs.UB);
+  Dir->setStrideVariable(Exprs.ST);
+  Dir->setEnsureUpperBound(Exprs.EUB);
+  Dir->setNextLowerBound(Exprs.NLB);
+  Dir->setNextUpperBound(Exprs.NUB);
+  Dir->setCounters(Exprs.Counters);
+  Dir->setPrivateCounters(Exprs.PrivateCounters);
+  Dir->setInits(Exprs.Inits);
+  Dir->setUpdates(Exprs.Updates);
+  Dir->setFinals(Exprs.Finals);
+  Dir->setHasCancel(HasCancel);
+  return Dir;
+}
+
+OMPTargetParallelForDirective *
+OMPTargetParallelForDirective::CreateEmpty(const ASTContext &C,
+                                           unsigned NumClauses,
+                                           unsigned CollapsedNum, EmptyShell) {
+  unsigned Size = llvm::alignTo(sizeof(OMPTargetParallelForDirective),
+                                llvm::alignOf<OMPClause *>());
+  void *Mem = C.Allocate(
+      Size + sizeof(OMPClause *) * NumClauses +
+      sizeof(Stmt *) * numLoopChildren(CollapsedNum, OMPD_target_parallel_for));
+  return new (Mem) OMPTargetParallelForDirective(CollapsedNum, NumClauses);
+}
+
 OMPTargetDataDirective *OMPTargetDataDirective::Create(
     const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
     ArrayRef<OMPClause *> Clauses, Stmt *AssociatedStmt) {

clang/lib/AST/StmtPrinter.cpp

@@ -1089,6 +1089,12 @@ void StmtPrinter::VisitOMPTargetParallelDirective(
   PrintOMPExecutableDirective(Node);
 }
 
+void StmtPrinter::VisitOMPTargetParallelForDirective(
+    OMPTargetParallelForDirective *Node) {
+  Indent() << "#pragma omp target parallel for ";
+  PrintOMPExecutableDirective(Node);
+}
+
 void StmtPrinter::VisitOMPTeamsDirective(OMPTeamsDirective *Node) {
   Indent() << "#pragma omp teams ";
   PrintOMPExecutableDirective(Node);

clang/lib/AST/StmtProfile.cpp

@@ -599,6 +599,11 @@ void StmtProfiler::VisitOMPTargetParallelDirective(
   VisitOMPExecutableDirective(S);
 }
 
+void StmtProfiler::VisitOMPTargetParallelForDirective(
+    const OMPTargetParallelForDirective *S) {
+  VisitOMPExecutableDirective(S);
+}
+
 void StmtProfiler::VisitOMPTeamsDirective(const OMPTeamsDirective *S) {
   VisitOMPExecutableDirective(S);
 }

clang/lib/Basic/OpenMPKinds.cpp

@@ -460,6 +460,16 @@ bool clang::isAllowedClauseForDirective(OpenMPDirectiveKind DKind,
 #define OPENMP_TARGET_PARALLEL_CLAUSE(Name)                                    \
   case OMPC_##Name:                                                            \
     return true;
+#include "clang/Basic/OpenMPKinds.def"
+    default:
+      break;
+    }
+    break;
+  case OMPD_target_parallel_for:
+    switch (CKind) {
+#define OPENMP_TARGET_PARALLEL_FOR_CLAUSE(Name)                                \
+  case OMPC_##Name:                                                            \
+    return true;
 #include "clang/Basic/OpenMPKinds.def"
     default:
       break;
@@ -552,17 +562,17 @@ bool clang::isAllowedClauseForDirective(OpenMPDirectiveKind DKind,
 bool clang::isOpenMPLoopDirective(OpenMPDirectiveKind DKind) {
   return DKind == OMPD_simd || DKind == OMPD_for || DKind == OMPD_for_simd ||
          DKind == OMPD_parallel_for || DKind == OMPD_parallel_for_simd ||
-         DKind == OMPD_taskloop ||
-         DKind == OMPD_taskloop_simd ||
-         DKind == OMPD_distribute; // TODO add next directives.
+         DKind == OMPD_taskloop || DKind == OMPD_taskloop_simd ||
+         DKind == OMPD_distribute ||
+         DKind == OMPD_target_parallel_for; // TODO add next directives.
 }
 
 bool clang::isOpenMPWorksharingDirective(OpenMPDirectiveKind DKind) {
   return DKind == OMPD_for || DKind == OMPD_for_simd ||
          DKind == OMPD_sections || DKind == OMPD_section ||
          DKind == OMPD_single || DKind == OMPD_parallel_for ||
-         DKind == OMPD_parallel_for_simd ||
-         DKind == OMPD_parallel_sections; // TODO add next directives.
+         DKind == OMPD_parallel_for_simd || DKind == OMPD_parallel_sections ||
+         DKind == OMPD_target_parallel_for; // TODO add next directives.
 }
 
 bool clang::isOpenMPTaskLoopDirective(OpenMPDirectiveKind DKind) {
@@ -571,14 +581,15 @@ bool clang::isOpenMPTaskLoopDirective(OpenMPDirectiveKind DKind) {
 
 bool clang::isOpenMPParallelDirective(OpenMPDirectiveKind DKind) {
   return DKind == OMPD_parallel || DKind == OMPD_parallel_for ||
-         DKind == OMPD_parallel_for_simd ||
-         DKind == OMPD_parallel_sections || DKind == OMPD_target_parallel;
-         // TODO add next directives.
+         DKind == OMPD_parallel_for_simd || DKind == OMPD_parallel_sections ||
+         DKind == OMPD_target_parallel || DKind == OMPD_target_parallel_for;
+  // TODO add next directives.
 }
 
 bool clang::isOpenMPTargetExecutionDirective(OpenMPDirectiveKind DKind) {
   // TODO add next directives.
-  return DKind == OMPD_target || DKind == OMPD_target_parallel;
+  return DKind == OMPD_target || DKind == OMPD_target_parallel ||
+         DKind == OMPD_target_parallel_for;
 }
 
 bool clang::isOpenMPTargetDataManagementDirective(OpenMPDirectiveKind DKind) {

clang/lib/CodeGen/CGStmt.cpp

@@ -265,6 +265,9 @@ void CodeGenFunction::EmitStmt(const Stmt *S) {
   case Stmt::OMPTargetParallelDirectiveClass:
     EmitOMPTargetParallelDirective(cast<OMPTargetParallelDirective>(*S));
     break;
+  case Stmt::OMPTargetParallelForDirectiveClass:
+    EmitOMPTargetParallelForDirective(cast<OMPTargetParallelForDirective>(*S));
+    break;
   case Stmt::OMPTaskLoopDirectiveClass:
     EmitOMPTaskLoopDirective(cast<OMPTaskLoopDirective>(*S));
     break;

clang/lib/CodeGen/CGStmtOpenMP.cpp

@@ -2715,6 +2715,11 @@ void CodeGenFunction::EmitOMPTargetParallelDirective(
   // TODO: codegen for target parallel.
 }
 
+void CodeGenFunction::EmitOMPTargetParallelForDirective(
+    const OMPTargetParallelForDirective &S) {
+  // TODO: codegen for target parallel for.
+}
+
 void CodeGenFunction::EmitOMPTaskLoopDirective(const OMPTaskLoopDirective &S) {
   // emit the code inside the construct for now
   auto CS = cast<CapturedStmt>(S.getAssociatedStmt());

clang/lib/CodeGen/CodeGenFunction.h

@@ -2344,6 +2344,8 @@ class CodeGenFunction : public CodeGenTypeCache {
   void EmitOMPTargetEnterDataDirective(const OMPTargetEnterDataDirective &S);
   void EmitOMPTargetExitDataDirective(const OMPTargetExitDataDirective &S);
   void EmitOMPTargetParallelDirective(const OMPTargetParallelDirective &S);
+  void
+  EmitOMPTargetParallelForDirective(const OMPTargetParallelForDirective &S);
   void EmitOMPTeamsDirective(const OMPTeamsDirective &S);
   void
   EmitOMPCancellationPointDirective(const OMPCancellationPointDirective &S);