[OpenACC] Implement AST/Sema for combined constructs

Combined constructs (OpenACC 3.3 section 2.11) are a short-cut for
writing a `loop` construct immediately inside of a `compute` construct.
However, this interaction requires we do additional work to ensure that
we get the semantics between the two correct, as well as diagnostics.

This patch adds the semantic analysis for the constructs (but no
    clauses), as well as the AST nodes.

Author: erichkeane

clang/include/clang-c/Index.h

@@ -2166,7 +2166,9 @@ enum CXCursorKind {
    */
   CXCursor_OpenACCLoopConstruct = 321,
 
-  CXCursor_LastStmt = CXCursor_OpenACCLoopConstruct,
+  CXCursor_OpenACCCombinedConstruct = 322,
+
+  CXCursor_LastStmt = CXCursor_OpenACCCombinedConstruct,
 
   /**
    * Cursor that represents the translation unit itself.

clang/include/clang/AST/RecursiveASTVisitor.h

@@ -4049,22 +4049,15 @@ bool RecursiveASTVisitor<Derived>::VisitOpenACCClauseList(
 
   for (const auto *C : Clauses)
     TRY_TO(VisitOpenACCClause(C));
-//    if (const auto *WithCond = dyn_cast<OopenACCClauseWithCondition>(C);
-//        WithCond && WIthCond->hasConditionExpr()) {
-//      TRY_TO(TraverseStmt(WithCond->getConditionExpr());
-//    } else if (const auto *
-//  }
-//  OpenACCClauseWithCondition::getConditionExpr/hasConditionExpr
-//OpenACCClauseWithExprs::children (might be null?)
-  // TODO OpenACC: When we have Clauses with expressions, we should visit them
-  // here.
   return true;
 }
 
 DEF_TRAVERSE_STMT(OpenACCComputeConstruct,
                   { TRY_TO(TraverseOpenACCAssociatedStmtConstruct(S)); })
 DEF_TRAVERSE_STMT(OpenACCLoopConstruct,
                   { TRY_TO(TraverseOpenACCAssociatedStmtConstruct(S)); })
+DEF_TRAVERSE_STMT(OpenACCCombinedConstruct,
+                  { TRY_TO(TraverseOpenACCAssociatedStmtConstruct(S)); })
 
 // Traverse HLSL: Out argument expression
 DEF_TRAVERSE_STMT(HLSLOutArgExpr, {})

clang/include/clang/AST/StmtOpenACC.h

@@ -166,11 +166,6 @@ class OpenACCComputeConstruct final
   }
 
   void setStructuredBlock(Stmt *S) { setAssociatedStmt(S); }
-  // Serialization helper function that searches the structured block for 'loop'
-  // constructs that should be associated with this, and sets their parent
-  // compute construct to this one. This isn't necessary normally, since we have
-  // the ability to record the state during parsing.
-  void findAndSetChildLoops();
 
 public:
   static bool classof(const Stmt *T) {
@@ -204,6 +199,8 @@ class OpenACCLoopConstruct final
   friend class ASTStmtWriter;
   friend class ASTStmtReader;
   friend class ASTContext;
+  friend class OpenACCAssociatedStmtConstruct;
+  friend class OpenACCCombinedConstruct;
   friend class OpenACCComputeConstruct;
 
   OpenACCLoopConstruct(unsigned NumClauses);
@@ -243,5 +240,57 @@ class OpenACCLoopConstruct final
     return ParentComputeConstructKind;
   }
 };
+
+// This class represents a 'combined' construct, which has a bunch of rules
+// shared with both loop and compute constructs.
+class OpenACCCombinedConstruct final
+    : public OpenACCAssociatedStmtConstruct,
+      public llvm::TrailingObjects<OpenACCCombinedConstruct,
+                                   const OpenACCClause *> {
+  OpenACCCombinedConstruct(unsigned NumClauses)
+      : OpenACCAssociatedStmtConstruct(
+            OpenACCCombinedConstructClass, OpenACCDirectiveKind::Invalid,
+            SourceLocation{}, SourceLocation{}, SourceLocation{},
+            /*AssociatedStmt=*/nullptr) {
+    std::uninitialized_value_construct(
+        getTrailingObjects<const OpenACCClause *>(),
+        getTrailingObjects<const OpenACCClause *>() + NumClauses);
+    setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
+                                  NumClauses));
+  }
+
+  OpenACCCombinedConstruct(OpenACCDirectiveKind K, SourceLocation Start,
+                           SourceLocation DirectiveLoc, SourceLocation End,
+                           ArrayRef<const OpenACCClause *> Clauses,
+                           Stmt *StructuredBlock)
+      : OpenACCAssociatedStmtConstruct(OpenACCCombinedConstructClass, K, Start,
+                                       DirectiveLoc, End, StructuredBlock) {
+    assert(isOpenACCCombinedDirectiveKind(K) &&
+           "Only parallel loop, serial loop, and kernels loop constructs "
+           "should be represented by this type");
+
+    std::uninitialized_copy(Clauses.begin(), Clauses.end(),
+                            getTrailingObjects<const OpenACCClause *>());
+    setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
+                                  Clauses.size()));
+  }
+  void setStructuredBlock(Stmt *S) { setAssociatedStmt(S); }
+
+public:
+  static bool classof(const Stmt *T) {
+    return T->getStmtClass() == OpenACCCombinedConstructClass;
+  }
+
+  static OpenACCCombinedConstruct *CreateEmpty(const ASTContext &C,
+                                               unsigned NumClauses);
+  static OpenACCCombinedConstruct *
+  Create(const ASTContext &C, OpenACCDirectiveKind K, SourceLocation Start,
+         SourceLocation DirectiveLoc, SourceLocation End,
+         ArrayRef<const OpenACCClause *> Clauses, Stmt *StructuredBlock);
+  Stmt *getLoop() { return getAssociatedStmt(); }
+  const Stmt *getLoop() const {
+    return const_cast<OpenACCCombinedConstruct *>(this)->getLoop();
+  }
+};
 } // namespace clang
 #endif // LLVM_CLANG_AST_STMTOPENACC_H

clang/include/clang/AST/TextNodeDumper.h

@@ -410,6 +410,7 @@ class TextNodeDumper
   void VisitHLSLOutArgExpr(const HLSLOutArgExpr *E);
   void VisitOpenACCConstructStmt(const OpenACCConstructStmt *S);
   void VisitOpenACCLoopConstruct(const OpenACCLoopConstruct *S);
+  void VisitOpenACCCombinedConstruct(const OpenACCCombinedConstruct *S);
   void VisitOpenACCAsteriskSizeExpr(const OpenACCAsteriskSizeExpr *S);
   void VisitEmbedExpr(const EmbedExpr *S);
   void VisitAtomicExpr(const AtomicExpr *AE);

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -12613,11 +12613,11 @@ def note_acc_previous_clause_here : Note<"previous clause is here">;
 def note_acc_previous_expr_here : Note<"previous expression is here">;
 def err_acc_branch_in_out_compute_construct
     : Error<"invalid %select{branch|return|throw}0 %select{out of|into}1 "
-            "OpenACC Compute Construct">;
+            "OpenACC Compute/Combined Construct">;
 def note_acc_branch_into_compute_construct
-    : Note<"invalid branch into OpenACC Compute Construct">;
+    : Note<"invalid branch into OpenACC Compute/Combined Construct">;
 def note_acc_branch_out_of_compute_construct
-    : Note<"invalid branch out of OpenACC Compute Construct">;
+    : Note<"invalid branch out of OpenACC Compute/Combined Construct">;
 def warn_acc_if_self_conflict
     : Warning<"OpenACC construct 'self' has no effect when an 'if' clause "
               "evaluates to true">,
@@ -12691,7 +12691,7 @@ def err_acc_reduction_composite_member_type :Error<
     "OpenACC 'reduction' composite variable must not have non-scalar field">;
 def note_acc_reduction_composite_member_loc : Note<"invalid field is here">;
 def err_acc_loop_not_for_loop
-    : Error<"OpenACC 'loop' construct can only be applied to a 'for' loop">;
+    : Error<"OpenACC '%0' construct can only be applied to a 'for' loop">;
 def note_acc_construct_here : Note<"'%0' construct is here">;
 def err_acc_loop_spec_conflict
     : Error<"OpenACC clause '%0' on '%1' construct conflicts with previous "
@@ -12749,16 +12749,16 @@ def err_reduction_op_mismatch
     : Error<"OpenACC 'reduction' variable must have the same operator in all "
             "nested constructs (%0 vs %1)">;
 def err_acc_loop_variable_type
-    : Error<"loop variable of loop associated with an OpenACC 'loop' construct "
+    : Error<"loop variable of loop associated with an OpenACC '%0' construct "
             "must be of integer, pointer, or random-access-iterator type (is "
-            "%0)">;
+            "%1)">;
 def err_acc_loop_variable
-    : Error<"OpenACC 'loop' construct must have initialization clause in "
+    : Error<"OpenACC '%0' construct must have initialization clause in "
             "canonical form ('var = init' or 'T var = init')">;
 def err_acc_loop_terminating_condition
-    : Error<"OpenACC 'loop' construct must have a terminating condition">;
+    : Error<"OpenACC '%0' construct must have a terminating condition">;
 def err_acc_loop_not_monotonic
-    : Error<"OpenACC 'loop' variable must monotonically increase or decrease "
+    : Error<"OpenACC '%0' variable must monotonically increase or decrease "
             "('++', '--', or compound assignment)">;
 
 // AMDGCN builtins diagnostics

clang/include/clang/Basic/OpenACCKinds.h

@@ -152,6 +152,12 @@ inline bool isOpenACCComputeDirectiveKind(OpenACCDirectiveKind K) {
          K == OpenACCDirectiveKind::Kernels;
 }
 
+inline bool isOpenACCCombinedDirectiveKind(OpenACCDirectiveKind K) {
+  return K == OpenACCDirectiveKind::ParallelLoop ||
+         K == OpenACCDirectiveKind::SerialLoop ||
+         K == OpenACCDirectiveKind::KernelsLoop;
+}
+
 enum class OpenACCAtomicKind : uint8_t {
   Read,
   Write,

clang/include/clang/Basic/StmtNodes.td

@@ -307,6 +307,7 @@ def OpenACCAssociatedStmtConstruct
     : StmtNode<OpenACCConstructStmt, /*abstract=*/1>;
 def OpenACCComputeConstruct : StmtNode<OpenACCAssociatedStmtConstruct>;
 def OpenACCLoopConstruct : StmtNode<OpenACCAssociatedStmtConstruct>;
+def OpenACCCombinedConstruct : StmtNode<OpenACCAssociatedStmtConstruct>;
 
 // OpenACC Additional Expressions.
 def OpenACCAsteriskSizeExpr : StmtNode<Expr>;

clang/include/clang/Sema/Scope.h

@@ -154,7 +154,8 @@ class Scope {
     /// depth of recursion.
     LambdaScope = 0x8000000,
     /// This is the scope of an OpenACC Compute Construct, which restricts
-    /// jumping into/out of it.
+    /// jumping into/out of it. We also use this to represent 'combined'
+    /// constructs, since they have the same behavior.
     OpenACCComputeConstructScope = 0x10000000,
 
     /// This is a scope of type alias declaration.

clang/include/clang/Sema/SemaOpenACC.h

@@ -170,10 +170,14 @@ class SemaOpenACC : public SemaBase {
   /// 'worker' clauses.
   SourceLocation LoopVectorClauseLoc;
   /// If there is a current 'active' loop construct that does NOT have a 'seq'
-  /// clause on it, this has that source location. This permits us to implement
-  /// the 'loop' restrictions on the loop variable. This can be extended via
-  /// 'collapse', so we need to keep this around for a while.
-  SourceLocation LoopWithoutSeqLoc;
+  /// clause on it, this has that source location and loop Directive 'kind'.
+  /// This permits us to implement the 'loop' restrictions on the loop variable.
+  /// This can be extended via 'collapse', so we need to keep this around for a
+  /// while.
+  struct LoopWithoutSeqCheckingInfo {
+    OpenACCDirectiveKind Kind = OpenACCDirectiveKind::Invalid;
+    SourceLocation Loc;
+  } LoopWithoutSeqInfo;
 
   // Redeclaration of the version in OpenACCClause.h.
   using DeviceTypeArgument = std::pair<IdentifierInfo *, SourceLocation>;
@@ -762,7 +766,7 @@ class SemaOpenACC : public SemaBase {
     SourceLocation OldLoopGangClauseOnKernelLoc;
     SourceLocation OldLoopWorkerClauseLoc;
     SourceLocation OldLoopVectorClauseLoc;
-    SourceLocation OldLoopWithoutSeqLoc;
+    LoopWithoutSeqCheckingInfo OldLoopWithoutSeqInfo;
     llvm::SmallVector<OpenACCReductionClause *> ActiveReductionClauses;
     LoopInConstructRAII LoopRAII;
 

clang/include/clang/Serialization/ASTBitCodes.h

@@ -2005,6 +2005,7 @@ enum StmtCode {
   // OpenACC Constructs/Exprs
   STMT_OPENACC_COMPUTE_CONSTRUCT,
   STMT_OPENACC_LOOP_CONSTRUCT,
+  STMT_OPENACC_COMBINED_CONSTRUCT,
   EXPR_OPENACC_ASTERISK_SIZE,
 
   // HLSL Constructs

clang/lib/AST/StmtOpenACC.cpp

@@ -89,3 +89,25 @@ OpenACCLoopConstruct *OpenACCLoopConstruct::Create(
       OpenACCLoopConstruct(ParentKind, BeginLoc, DirLoc, EndLoc, Clauses, Loop);
   return Inst;
 }
+
+OpenACCCombinedConstruct *
+OpenACCCombinedConstruct::CreateEmpty(const ASTContext &C,
+                                      unsigned NumClauses) {
+  void *Mem = C.Allocate(
+      OpenACCCombinedConstruct::totalSizeToAlloc<const OpenACCClause *>(
+          NumClauses));
+  auto *Inst = new (Mem) OpenACCCombinedConstruct(NumClauses);
+  return Inst;
+}
+
+OpenACCCombinedConstruct *OpenACCCombinedConstruct::Create(
+    const ASTContext &C, OpenACCDirectiveKind DK, SourceLocation BeginLoc,
+    SourceLocation DirLoc, SourceLocation EndLoc,
+    ArrayRef<const OpenACCClause *> Clauses, Stmt *Loop) {
+  void *Mem = C.Allocate(
+      OpenACCCombinedConstruct::totalSizeToAlloc<const OpenACCClause *>(
+          Clauses.size()));
+  auto *Inst = new (Mem)
+      OpenACCCombinedConstruct(DK, BeginLoc, DirLoc, EndLoc, Clauses, Loop);
+  return Inst;
+}

clang/lib/AST/StmtPrinter.cpp

@@ -1184,6 +1184,18 @@ void StmtPrinter::VisitOpenACCLoopConstruct(OpenACCLoopConstruct *S) {
   PrintStmt(S->getLoop());
 }
 
+void StmtPrinter::VisitOpenACCCombinedConstruct(OpenACCCombinedConstruct *S) {
+  Indent() << "#pragma acc " << S->getDirectiveKind();
+  if (!S->clauses().empty()) {
+    OS << ' ';
+    OpenACCClausePrinter Printer(OS, Policy);
+    Printer.VisitClauseList(S->clauses());
+  }
+  OS << '\n';
+
+  PrintStmt(S->getLoop());
+}
+
 //===----------------------------------------------------------------------===//
 //  Expr printing methods.
 //===----------------------------------------------------------------------===//

clang/lib/AST/StmtProfile.cpp

@@ -2688,6 +2688,15 @@ void StmtProfiler::VisitOpenACCLoopConstruct(const OpenACCLoopConstruct *S) {
   P.VisitOpenACCClauseList(S->clauses());
 }
 
+void StmtProfiler::VisitOpenACCCombinedConstruct(
+    const OpenACCCombinedConstruct *S) {
+  // VisitStmt handles children, so the Loop is handled.
+  VisitStmt(S);
+
+  OpenACCClauseProfiler P{*this};
+  P.VisitOpenACCClauseList(S->clauses());
+}
+
 void StmtProfiler::VisitHLSLOutArgExpr(const HLSLOutArgExpr *S) {
   VisitStmt(S);
 }

clang/lib/AST/TextNodeDumper.cpp

@@ -2931,6 +2931,11 @@ void TextNodeDumper::VisitOpenACCLoopConstruct(const OpenACCLoopConstruct *S) {
     OS << " parent: " << S->getParentComputeConstructKind();
 }
 
+void TextNodeDumper::VisitOpenACCCombinedConstruct(
+    const OpenACCCombinedConstruct *S) {
+  OS << " " << S->getDirectiveKind();
+}
+
 void TextNodeDumper::VisitEmbedExpr(const EmbedExpr *S) {
   AddChild("begin", [=] { OS << S->getStartingElementPos(); });
   AddChild("number of elements", [=] { OS << S->getDataElementCount(); });

clang/lib/CodeGen/CGStmt.cpp

@@ -455,6 +455,9 @@ void CodeGenFunction::EmitStmt(const Stmt *S, ArrayRef<const Attr *> Attrs) {
   case Stmt::OpenACCLoopConstructClass:
     EmitOpenACCLoopConstruct(cast<OpenACCLoopConstruct>(*S));
     break;
+  case Stmt::OpenACCCombinedConstructClass:
+    EmitOpenACCCombinedConstruct(cast<OpenACCCombinedConstruct>(*S));
+    break;
   }
 }
 

clang/lib/CodeGen/CodeGenFunction.h

@@ -4087,6 +4087,13 @@ class CodeGenFunction : public CodeGenTypeCache {
     EmitStmt(S.getLoop());
   }
 
+  void EmitOpenACCCombinedConstruct(const OpenACCCombinedConstruct &S) {
+    // TODO OpenACC: Implement this.  It is currently implemented as a 'no-op',
+    // simply emitting its loop, but in the future we will implement
+    // some sort of IR.
+    EmitStmt(S.getLoop());
+  }
+
   //===--------------------------------------------------------------------===//
   //                         LValue Expression Emission
   //===--------------------------------------------------------------------===//

clang/lib/Parse/ParseOpenACC.cpp

@@ -571,6 +571,9 @@ bool doesDirectiveHaveAssociatedStmt(OpenACCDirectiveKind DirKind) {
   case OpenACCDirectiveKind::Parallel:
   case OpenACCDirectiveKind::Serial:
   case OpenACCDirectiveKind::Kernels:
+  case OpenACCDirectiveKind::ParallelLoop:
+  case OpenACCDirectiveKind::SerialLoop:
+  case OpenACCDirectiveKind::KernelsLoop:
   case OpenACCDirectiveKind::Loop:
     return true;
   }
@@ -582,6 +585,9 @@ unsigned getOpenACCScopeFlags(OpenACCDirectiveKind DirKind) {
   case OpenACCDirectiveKind::Parallel:
   case OpenACCDirectiveKind::Serial:
   case OpenACCDirectiveKind::Kernels:
+  case OpenACCDirectiveKind::ParallelLoop:
+  case OpenACCDirectiveKind::SerialLoop:
+  case OpenACCDirectiveKind::KernelsLoop:
     // Mark this as a BreakScope/ContinueScope as well as a compute construct
     // so that we can diagnose trying to 'break'/'continue' inside of one.
     return Scope::BreakScope | Scope::ContinueScope |

clang/lib/Sema/JumpDiagnostics.cpp

@@ -616,6 +616,16 @@ void JumpScopeChecker::BuildScopeInformation(Stmt *S,
     return;
   }
 
+  case Stmt::OpenACCCombinedConstructClass: {
+    unsigned NewParentScope = Scopes.size();
+    OpenACCCombinedConstruct *CC = cast<OpenACCCombinedConstruct>(S);
+    Scopes.push_back(GotoScope(
+        ParentScope, diag::note_acc_branch_into_compute_construct,
+        diag::note_acc_branch_out_of_compute_construct, CC->getBeginLoc()));
+    BuildScopeInformation(CC->getLoop(), NewParentScope);
+    return;
+  }
+
   default:
     if (auto *ED = dyn_cast<OMPExecutableDirective>(S)) {
       if (!ED->isStandaloneDirective()) {

clang/lib/Sema/SemaExceptionSpec.cpp

@@ -1406,6 +1406,7 @@ CanThrowResult Sema::canThrow(const Stmt *S) {
     // Most statements can throw if any substatement can throw.
   case Stmt::OpenACCComputeConstructClass:
   case Stmt::OpenACCLoopConstructClass:
+  case Stmt::OpenACCCombinedConstructClass:
   case Stmt::AttributedStmtClass:
   case Stmt::BreakStmtClass:
   case Stmt::CapturedStmtClass: