Allow multiple tile clauses, fix multiple 'collapse' checking

There is no restriction on multiple tile clauses.  There is perhaps some
goofy behavioral decisions to make during codegen to decide which one
wins, but here it isn't too bad.

Multiple 'collapse' clauses have a similar problem, but with more
difficult checking (since we have to take the instantiation status into
    consideration).

Author: erichkeane

clang/include/clang/Sema/SemaOpenACC.h

@@ -81,7 +81,7 @@ class SemaOpenACC : public SemaBase {
   /// Collapse has an 'N' count that makes it apply to a number of loops 'below'
   /// it.
   struct CollapseCheckingInfo {
-    OpenACCCollapseClause *ActiveCollapse = nullptr;
+    const OpenACCCollapseClause *ActiveCollapse = nullptr;
 
     /// This is a value that maintains the current value of the 'N' on the
     /// current collapse, minus the depth that has already been traversed. When
@@ -110,7 +110,7 @@ class SemaOpenACC : public SemaBase {
     /// own counting of elements.
     UnsignedOrNone CurTileCount = std::nullopt;
 
-    /// Records whether we've hit a 'CurTileCount' of '0' on the wya down,
+    /// Records whether we've hit a 'CurTileCount' of '0' on the way down,
     /// which allows us to diagnose if the number of arguments is too large for
     /// the current number of 'for' loops.
     bool TileDepthSatisfied = true;

clang/lib/AST/OpenACCClause.cpp

@@ -190,18 +190,15 @@ OpenACCCollapseClause::OpenACCCollapseClause(SourceLocation BeginLoc,
                                              SourceLocation EndLoc)
     : OpenACCClauseWithSingleIntExpr(OpenACCClauseKind::Collapse, BeginLoc,
                                      LParenLoc, LoopCount, EndLoc),
-      HasForce(HasForce) {
-  assert(LoopCount && "LoopCount required");
-}
+      HasForce(HasForce) {}
 
 OpenACCCollapseClause *
 OpenACCCollapseClause::Create(const ASTContext &C, SourceLocation BeginLoc,
                               SourceLocation LParenLoc, bool HasForce,
                               Expr *LoopCount, SourceLocation EndLoc) {
-  assert(
-      LoopCount &&
-      (LoopCount->isInstantiationDependent() || isa<ConstantExpr>(LoopCount)) &&
-      "Loop count not constant expression");
+  assert((!LoopCount || (LoopCount->isInstantiationDependent() ||
+                         isa<ConstantExpr>(LoopCount))) &&
+         "Loop count not constant expression");
   void *Mem =
       C.Allocate(sizeof(OpenACCCollapseClause), alignof(OpenACCCollapseClause));
   return new (Mem)

clang/lib/Sema/SemaOpenACC.cpp

@@ -221,50 +221,106 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
   }
 }
 
+namespace {
+// Given two collapse clauses, and the uninstanted version of the new one,
+// return the 'best' one for the purposes of setting the collapse checking
+// values.
+const OpenACCCollapseClause *
+getBestCollapseCandidate(const OpenACCCollapseClause *Old,
+                         const OpenACCCollapseClause *New,
+                         const OpenACCCollapseClause *UnInstNew) {
+  // If the loop count is nullptr, it is because instantiation failed, so this
+  // can't be the best one.
+  if (!New->getLoopCount())
+    return Old;
+
+  // If the loop-count had an error, than 'new' isn't a candidate.
+  if (!New->getLoopCount())
+    return Old;
+
+  // Don't consider uninstantiated ones, since we can't really check these.
+  if (New->getLoopCount()->isInstantiationDependent())
+    return Old;
+
+  // If this is an instantiation, and the old version wasn't instantation
+  // dependent, than nothing has changed and we've already done a diagnostic
+  // based on this one, so don't consider it.
+  if (UnInstNew && !UnInstNew->getLoopCount()->isInstantiationDependent())
+    return Old;
+
+  // New is now a valid candidate, so if there isn't an old one at this point,
+  // New is the only valid one.
+  if (!Old)
+    return New;
+
+  // If the 'New' expression has a larger value than 'Old', then it is the new
+  // best candidate.
+  if (cast<ConstantExpr>(Old->getLoopCount())->getResultAsAPSInt() <
+      cast<ConstantExpr>(New->getLoopCount())->getResultAsAPSInt())
+    return New;
+
+  return Old;
+}
+} // namespace
+
 void SemaOpenACC::AssociatedStmtRAII::SetCollapseInfoBeforeAssociatedStmt(
     ArrayRef<const OpenACCClause *> UnInstClauses,
     ArrayRef<OpenACCClause *> Clauses) {
 
   // Reset this checking for loops that aren't covered in a RAII object.
   SemaRef.LoopInfo.CurLevelHasLoopAlready = false;
   SemaRef.CollapseInfo.CollapseDepthSatisfied = true;
+  SemaRef.CollapseInfo.CurCollapseCount = 0;
   SemaRef.TileInfo.TileDepthSatisfied = true;
 
   // We make sure to take an optional list of uninstantiated clauses, so that
   // we can check to make sure we don't 'double diagnose' in the event that
-  // the value of 'N' was not dependent in a template. We also ensure during
-  // Sema that there is only 1 collapse on each construct, so we can count on
-  // the fact that if both find a 'collapse', that they are the same one.
-  auto *CollapseClauseItr =
-      llvm::find_if(Clauses, llvm::IsaPred<OpenACCCollapseClause>);
-  auto *UnInstCollapseClauseItr =
+  // the value of 'N' was not dependent in a template. Since we cannot count on
+  // there only being a single collapse clause, we count on the order to make
+  // sure get the matching ones, and we count on TreeTransform not removing
+  // these, even if loop-count instantiation failed. We can check the
+  // non-dependent ones right away, and realize that subsequent instantiation
+  // can only make it more specific.
+
+  auto *UnInstClauseItr =
       llvm::find_if(UnInstClauses, llvm::IsaPred<OpenACCCollapseClause>);
-
-  if (Clauses.end() == CollapseClauseItr)
-    return;
-
-  OpenACCCollapseClause *CollapseClause =
-      cast<OpenACCCollapseClause>(*CollapseClauseItr);
-
-  SemaRef.CollapseInfo.ActiveCollapse = CollapseClause;
-  Expr *LoopCount = CollapseClause->getLoopCount();
-
-  // If the loop count is still instantiation dependent, setting the depth
-  // counter isn't necessary, so return here.
-  if (!LoopCount || LoopCount->isInstantiationDependent())
-    return;
-
-  // Suppress diagnostics if we've done a 'transform' where the previous version
-  // wasn't dependent, meaning we already diagnosed it.
-  if (UnInstCollapseClauseItr != UnInstClauses.end() &&
-      !cast<OpenACCCollapseClause>(*UnInstCollapseClauseItr)
-           ->getLoopCount()
-           ->isInstantiationDependent())
+  auto *ClauseItr =
+      llvm::find_if(Clauses, llvm::IsaPred<OpenACCCollapseClause>);
+  const OpenACCCollapseClause *FoundClause = nullptr;
+
+  // Loop through the list of Collapse clauses and find the one that:
+  // 1- Has a non-dependent, non-null loop count (null means error, likely
+  // during instantiation).
+  // 2- If UnInstClauses isn't empty, its corresponding
+  // loop count was dependent.
+  // 3- Has the largest 'loop count' of all.
+  while (ClauseItr != Clauses.end()) {
+    const OpenACCCollapseClause *CurClause =
+        cast<OpenACCCollapseClause>(*ClauseItr);
+    const OpenACCCollapseClause *UnInstCurClause =
+        UnInstClauseItr == UnInstClauses.end()
+            ? nullptr
+            : cast<OpenACCCollapseClause>(*UnInstClauseItr);
+
+    FoundClause =
+        getBestCollapseCandidate(FoundClause, CurClause, UnInstCurClause);
+
+    UnInstClauseItr =
+        UnInstClauseItr == UnInstClauses.end()
+            ? UnInstClauseItr
+            : std::find_if(std::next(UnInstClauseItr), UnInstClauses.end(),
+                           llvm::IsaPred<OpenACCCollapseClause>);
+    ClauseItr = std::find_if(std::next(ClauseItr), Clauses.end(),
+                             llvm::IsaPred<OpenACCCollapseClause>);
+  }
+
+  if (!FoundClause)
     return;
 
+  SemaRef.CollapseInfo.ActiveCollapse = FoundClause;
   SemaRef.CollapseInfo.CollapseDepthSatisfied = false;
   SemaRef.CollapseInfo.CurCollapseCount =
-      cast<ConstantExpr>(LoopCount)->getResultAsAPSInt();
+      cast<ConstantExpr>(FoundClause->getLoopCount())->getResultAsAPSInt();
   SemaRef.CollapseInfo.DirectiveKind = DirKind;
 }
 
@@ -283,6 +339,17 @@ void SemaOpenACC::AssociatedStmtRAII::SetTileInfoBeforeAssociatedStmt(
     return;
 
   OpenACCTileClause *TileClause = cast<OpenACCTileClause>(*TileClauseItr);
+
+  // Multiple tile clauses are allowed, so ensure that we use the one with the
+  // largest 'tile count'.
+  while (Clauses.end() !=
+         (TileClauseItr = std::find_if(std::next(TileClauseItr), Clauses.end(),
+                                       llvm::IsaPred<OpenACCTileClause>))) {
+    OpenACCTileClause *NewClause = cast<OpenACCTileClause>(*TileClauseItr);
+    if (NewClause->getSizeExprs().size() > TileClause->getSizeExprs().size())
+      TileClause = NewClause;
+  }
+
   SemaRef.TileInfo.ActiveTile = TileClause;
   SemaRef.TileInfo.TileDepthSatisfied = false;
   SemaRef.TileInfo.CurTileCount =

clang/lib/Sema/TreeTransform.h

@@ -12344,9 +12344,6 @@ void OpenACCClauseTransform<Derived>::VisitCollapseClause(
   NewLoopCount =
       Self.getSema().OpenACC().CheckCollapseLoopCount(NewLoopCount.get());
 
-  if (!NewLoopCount.isUsable())
-    return;
-
   ParsedClause.setCollapseDetails(C.hasForce(), NewLoopCount.get());
   NewClause = OpenACCCollapseClause::Create(
       Self.getSema().getASTContext(), ParsedClause.getBeginLoc(),

clang/test/SemaOpenACC/combined-construct-tile-clause.cpp

@@ -143,13 +143,6 @@ void only_for_loops() {
     do{}while(true);
 }
 
-void only_one_on_loop() {
-  // expected-error@+2{{OpenACC 'tile' clause cannot appear more than once on a 'serial loop' directive}}
-  // expected-note@+1{{previous clause is here}}
-#pragma acc serial loop tile(1) tile(1)
-  for(int i = 0; i < 5; ++i);
-}
-
 template<unsigned Val>
 void depth_too_high_templ() {
   // expected-error@+1{{'tile' clause specifies a loop count greater than the number of available loops}}

clang/test/SemaOpenACC/loop-construct-collapse-clause.cpp

@@ -493,3 +493,30 @@ void intervening_without_force() {
       for(;;);
 }
 
+template<unsigned N>
+void allow_multiple_collapse_templ() {
+    // expected-error@+2{{'collapse' clause specifies a loop count greater than the number of available loops}}
+    // expected-note@+1{{active 'collapse' clause defined here}}
+#pragma acc loop collapse(N) collapse(N+1)
+  for(unsigned i = 0; i < 5; ++i)
+    for(unsigned j = 0; j < 5; ++j);
+    // expected-error@+2{{'collapse' clause specifies a loop count greater than the number of available loops}}
+    // expected-note@+1{{active 'collapse' clause defined here}}
+#pragma acc loop collapse(N+1) collapse(N)
+  for(unsigned i = 0; i < 5; ++i)
+    for(unsigned j = 0; j < 5; ++j);
+}
+void allow_multiple_collapse() {
+  allow_multiple_collapse_templ<2>(); // expected-note{{in instantiation}}
+
+    // expected-error@+2{{'collapse' clause specifies a loop count greater than the number of available loops}}
+    // expected-note@+1{{active 'collapse' clause defined here}}
+#pragma acc loop collapse(2) collapse(3)
+  for(unsigned i = 0; i < 5; ++i)
+    for(unsigned j = 0; j < 5; ++j);
+    // expected-error@+2{{'collapse' clause specifies a loop count greater than the number of available loops}}
+    // expected-note@+1{{active 'collapse' clause defined here}}
+#pragma acc loop collapse(3) collapse(2)
+  for(unsigned i = 0; i < 5; ++i)
+    for(unsigned j = 0; j < 5; ++j);
+}

clang/test/SemaOpenACC/loop-construct-tile-clause.cpp

@@ -143,13 +143,6 @@ void only_for_loops() {
     do{}while(true);
 }
 
-void only_one_on_loop() {
-  // expected-error@+2{{OpenACC 'tile' clause cannot appear more than once on a 'loop' directive}}
-  // expected-note@+1{{previous clause is here}}
-#pragma acc loop tile(1) tile(1)
-  for(int i = 0; i < 5; ++i);
-}
-
 template<unsigned Val>
 void depth_too_high_templ() {
   // expected-error@+1{{'tile' clause specifies a loop count greater than the number of available loops}}
@@ -384,6 +377,20 @@ void intervening() {
   }
 }
 
+void use_largest_tile() {
+// expected-error@+2{{'tile' clause specifies a loop count greater than the number of available loops}}
+// expected-note@+1{{active 'tile' clause defined here}}
+#pragma acc loop tile(1,2) tile (3,4,5)
+  for(int i = 0; i < 5; ++i)
+    for (int j = 0; j < 5; ++j);
+
+// expected-error@+2{{'tile' clause specifies a loop count greater than the number of available loops}}
+// expected-note@+1{{active 'tile' clause defined here}}
+#pragma acc loop tile (3,4,5) tile(1,2)
+  for(int i = 0; i < 5; ++i)
+    for (int j = 0; j < 5; ++j);
+}
+
 void collapse_tile_depth() {
   // expected-error@+4{{'collapse' clause specifies a loop count greater than the number of available loops}}
   // expected-note@+3{{active 'collapse' clause defined here}}