[OpenACC] Implement 'tile' loop count/tightly nested loop requirement (#111038)

the 'tile' clause requires that it be followed by N (where N is the
    number of size expressions) 'tightly nested loops'.  This means the
same as it does in 'collapse', so much of the implementation is
simliar/shared with that.

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -12670,19 +12670,19 @@ def err_acc_size_expr_value
     : Error<
           "OpenACC 'tile' clause size expression must be %select{an asterisk "
           "or a constant expression|positive integer value, evaluated to %1}0">;
-def err_acc_invalid_in_collapse_loop
-    : Error<"%select{OpenACC '%1' construct|while loop|do loop}0 cannot appear "
-            "in intervening code of a 'loop' with a 'collapse' clause">;
-def note_acc_collapse_clause_here
-    : Note<"active 'collapse' clause defined here">;
-def err_acc_collapse_multiple_loops
+def err_acc_invalid_in_loop
+    : Error<"%select{OpenACC '%2' construct|while loop|do loop}0 cannot appear "
+            "in intervening code of a 'loop' with a '%1' clause">;
+def note_acc_active_clause_here
+    : Note<"active '%0' clause defined here">;
+def err_acc_clause_multiple_loops
     : Error<"more than one for-loop in a loop associated with OpenACC 'loop' "
-            "construct with a 'collapse' clause">;
-def err_acc_collapse_insufficient_loops
-    : Error<"'collapse' clause specifies a loop count greater than the number "
+            "construct with a '%select{collapse|tile}0' clause">;
+def err_acc_insufficient_loops
+    : Error<"'%0' clause specifies a loop count greater than the number "
             "of available loops">;
-def err_acc_collapse_intervening_code
-    : Error<"inner loops must be tightly nested inside a 'collapse' clause on "
+def err_acc_intervening_code
+    : Error<"inner loops must be tightly nested inside a '%0' clause on "
             "a 'loop' construct">;
 
 // AMDGCN builtins diagnostics

clang/include/clang/Sema/SemaOpenACC.h

@@ -42,6 +42,23 @@ private:
   /// above collection.
   bool InsideComputeConstruct = false;
 
+  /// Certain clauses care about the same things that aren't specific to the
+  /// individual clause, but can be shared by a few, so store them here. All
+  /// require a 'no intervening constructs' rule, so we know they are all from
+  /// the same 'place'.
+  struct LoopCheckingInfo {
+    /// Records whether we've seen the top level 'for'. We already diagnose
+    /// later that the 'top level' is a for loop, so we use this to suppress the
+    /// 'collapse inner loop not a 'for' loop' diagnostic.
+    LLVM_PREFERRED_TYPE(bool)
+    unsigned TopLevelLoopSeen : 1;
+
+    /// Records whether this 'tier' of the loop has already seen a 'for' loop,
+    /// used to diagnose if there are multiple 'for' loops at any one level.
+    LLVM_PREFERRED_TYPE(bool)
+    unsigned CurLevelHasLoopAlready : 1;
+  } LoopInfo{/*TopLevelLoopSeen=*/false, /*CurLevelHasLoopAlready=*/false};
+
   /// The 'collapse' clause requires quite a bit of checking while
   /// parsing/instantiating its body, so this structure/object keeps all of the
   /// necessary information as we do checking.  This should rarely be directly
@@ -59,25 +76,27 @@ private:
     /// else it should be 'N' minus the current depth traversed.
     std::optional<llvm::APSInt> CurCollapseCount;
 
-    /// Records whether we've seen the top level 'for'. We already diagnose
-    /// later that the 'top level' is a for loop, so we use this to suppress the
-    /// 'collapse inner loop not a 'for' loop' diagnostic.
-    LLVM_PREFERRED_TYPE(bool)
-    unsigned TopLevelLoopSeen : 1;
-
-    /// Records whether this 'tier' of the loop has already seen a 'for' loop,
-    /// used to diagnose if there are multiple 'for' loops at any one level.
-    LLVM_PREFERRED_TYPE(bool)
-    unsigned CurLevelHasLoopAlready : 1;
-
     /// Records whether we've hit a CurCollapseCount of '0' on the way down,
     /// which allows us to diagnose if the value of 'N' is too large for the
     /// current number of 'for' loops.
-    LLVM_PREFERRED_TYPE(bool)
-    unsigned CollapseDepthSatisfied : 1;
-  } CollapseInfo{nullptr, std::nullopt, /*TopLevelLoopSeen=*/false,
-                 /*CurLevelHasLoopAlready=*/false,
-                 /*CollapseDepthSatisfied=*/true};
+    bool CollapseDepthSatisfied = true;
+  } CollapseInfo;
+
+  /// The 'tile' clause requires a bit of additional checking as well, so like
+  /// the `CollapseCheckingInfo`, ensure we maintain information here too.
+  struct TileCheckingInfo {
+    OpenACCTileClause *ActiveTile = nullptr;
+
+    /// This is the number of expressions on a 'tile' clause.  This doesn't have
+    /// to be an APSInt because it isn't the result of a constexpr, just by our
+    /// own counting of elements.
+    std::optional<unsigned> CurTileCount;
+
+    /// Records whether we've hit a 'CurTileCount' of '0' on the wya down,
+    /// which allows us to diagnose if the number of arguments is too large for
+    /// the current number of 'for' loops.
+    bool TileDepthSatisfied = true;
+  } TileInfo;
 
 public:
   // Redeclaration of the version in OpenACCClause.h.
@@ -537,12 +556,15 @@ public:
   /// into a loop (for, etc) inside the construct.
   class LoopInConstructRAII {
     SemaOpenACC &SemaRef;
+    LoopCheckingInfo OldLoopInfo;
     CollapseCheckingInfo OldCollapseInfo;
+    TileCheckingInfo OldTileInfo;
     bool PreserveDepth;
 
   public:
     LoopInConstructRAII(SemaOpenACC &SemaRef, bool PreserveDepth = true)
-        : SemaRef(SemaRef), OldCollapseInfo(SemaRef.CollapseInfo),
+        : SemaRef(SemaRef), OldLoopInfo(SemaRef.LoopInfo),
+          OldCollapseInfo(SemaRef.CollapseInfo), OldTileInfo(SemaRef.TileInfo),
           PreserveDepth(PreserveDepth) {}
     ~LoopInConstructRAII() {
       // The associated-statement level of this should NOT preserve this, as it
@@ -551,12 +573,20 @@ public:
       bool CollapseDepthSatisified =
           PreserveDepth ? SemaRef.CollapseInfo.CollapseDepthSatisfied
                         : OldCollapseInfo.CollapseDepthSatisfied;
+      bool TileDepthSatisfied = PreserveDepth
+                                    ? SemaRef.TileInfo.TileDepthSatisfied
+                                    : OldTileInfo.TileDepthSatisfied;
       bool CurLevelHasLoopAlready =
-          PreserveDepth ? SemaRef.CollapseInfo.CurLevelHasLoopAlready
-                        : OldCollapseInfo.CurLevelHasLoopAlready;
+          PreserveDepth ? SemaRef.LoopInfo.CurLevelHasLoopAlready
+                        : OldLoopInfo.CurLevelHasLoopAlready;
+
+      SemaRef.LoopInfo = OldLoopInfo;
       SemaRef.CollapseInfo = OldCollapseInfo;
+      SemaRef.TileInfo = OldTileInfo;
+
       SemaRef.CollapseInfo.CollapseDepthSatisfied = CollapseDepthSatisified;
-      SemaRef.CollapseInfo.CurLevelHasLoopAlready = CurLevelHasLoopAlready;
+      SemaRef.TileInfo.TileDepthSatisfied = TileDepthSatisfied;
+      SemaRef.LoopInfo.CurLevelHasLoopAlready = CurLevelHasLoopAlready;
     }
   };
 
@@ -577,6 +607,9 @@ public:
     void SetCollapseInfoBeforeAssociatedStmt(
         ArrayRef<const OpenACCClause *> UnInstClauses,
         ArrayRef<OpenACCClause *> Clauses);
+    void SetTileInfoBeforeAssociatedStmt(
+        ArrayRef<const OpenACCClause *> UnInstClauses,
+        ArrayRef<OpenACCClause *> Clauses);
     ~AssociatedStmtRAII();
   };
 };

clang/lib/Sema/SemaOpenACC.cpp

@@ -1128,6 +1128,7 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
     SemaRef.ParentlessLoopConstructs.swap(ParentlessLoopConstructs);
   } else if (DirKind == OpenACCDirectiveKind::Loop) {
     SetCollapseInfoBeforeAssociatedStmt(UnInstClauses, Clauses);
+    SetTileInfoBeforeAssociatedStmt(UnInstClauses, Clauses);
   }
 }
 
@@ -1136,8 +1137,9 @@ void SemaOpenACC::AssociatedStmtRAII::SetCollapseInfoBeforeAssociatedStmt(
     ArrayRef<OpenACCClause *> Clauses) {
 
   // Reset this checking for loops that aren't covered in a RAII object.
-  SemaRef.CollapseInfo.CurLevelHasLoopAlready = false;
+  SemaRef.LoopInfo.CurLevelHasLoopAlready = false;
   SemaRef.CollapseInfo.CollapseDepthSatisfied = true;
+  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
@@ -1176,6 +1178,26 @@ void SemaOpenACC::AssociatedStmtRAII::SetCollapseInfoBeforeAssociatedStmt(
       cast<ConstantExpr>(LoopCount)->getResultAsAPSInt();
 }
 
+void SemaOpenACC::AssociatedStmtRAII::SetTileInfoBeforeAssociatedStmt(
+    ArrayRef<const OpenACCClause *> UnInstClauses,
+    ArrayRef<OpenACCClause *> Clauses) {
+  // We don't diagnose if this is during instantiation, since the only thing we
+  // care about is the number of arguments, which we can figure out without
+  // instantiation, so we don't want to double-diagnose.
+  if (UnInstClauses.size() > 0)
+    return;
+  auto *TileClauseItr =
+      llvm::find_if(Clauses, llvm::IsaPred<OpenACCTileClause>);
+
+  if (Clauses.end() == TileClauseItr)
+    return;
+
+  OpenACCTileClause *TileClause = cast<OpenACCTileClause>(*TileClauseItr);
+  SemaRef.TileInfo.ActiveTile = TileClause;
+  SemaRef.TileInfo.TileDepthSatisfied = false;
+  SemaRef.TileInfo.CurTileCount = TileClause->getSizeExprs().size();
+}
+
 SemaOpenACC::AssociatedStmtRAII::~AssociatedStmtRAII() {
   SemaRef.InsideComputeConstruct = WasInsideComputeConstruct;
   if (DirKind == OpenACCDirectiveKind::Parallel ||
@@ -1771,38 +1793,66 @@ void SemaOpenACC::ActOnWhileStmt(SourceLocation WhileLoc) {
   if (!getLangOpts().OpenACC)
     return;
 
-  if (!CollapseInfo.TopLevelLoopSeen)
+  if (!LoopInfo.TopLevelLoopSeen)
     return;
 
   if (CollapseInfo.CurCollapseCount && *CollapseInfo.CurCollapseCount > 0) {
-    Diag(WhileLoc, diag::err_acc_invalid_in_collapse_loop) << /*while loop*/ 1;
+    Diag(WhileLoc, diag::err_acc_invalid_in_loop)
+        << /*while loop*/ 1 << OpenACCClauseKind::Collapse;
     assert(CollapseInfo.ActiveCollapse && "Collapse count without object?");
     Diag(CollapseInfo.ActiveCollapse->getBeginLoc(),
-         diag::note_acc_collapse_clause_here);
+         diag::note_acc_active_clause_here)
+        << OpenACCClauseKind::Collapse;
 
     // Remove the value so that we don't get cascading errors in the body. The
     // caller RAII object will restore this.
     CollapseInfo.CurCollapseCount = std::nullopt;
   }
+
+  if (TileInfo.CurTileCount && *TileInfo.CurTileCount > 0) {
+    Diag(WhileLoc, diag::err_acc_invalid_in_loop)
+        << /*while loop*/ 1 << OpenACCClauseKind::Tile;
+    assert(TileInfo.ActiveTile && "tile count without object?");
+    Diag(TileInfo.ActiveTile->getBeginLoc(), diag::note_acc_active_clause_here)
+        << OpenACCClauseKind::Tile;
+
+    // Remove the value so that we don't get cascading errors in the body. The
+    // caller RAII object will restore this.
+    TileInfo.CurTileCount = std::nullopt;
+  }
 }
 
 void SemaOpenACC::ActOnDoStmt(SourceLocation DoLoc) {
   if (!getLangOpts().OpenACC)
     return;
 
-  if (!CollapseInfo.TopLevelLoopSeen)
+  if (!LoopInfo.TopLevelLoopSeen)
     return;
 
   if (CollapseInfo.CurCollapseCount && *CollapseInfo.CurCollapseCount > 0) {
-    Diag(DoLoc, diag::err_acc_invalid_in_collapse_loop) << /*do loop*/ 2;
+    Diag(DoLoc, diag::err_acc_invalid_in_loop)
+        << /*do loop*/ 2 << OpenACCClauseKind::Collapse;
     assert(CollapseInfo.ActiveCollapse && "Collapse count without object?");
     Diag(CollapseInfo.ActiveCollapse->getBeginLoc(),
-         diag::note_acc_collapse_clause_here);
+         diag::note_acc_active_clause_here)
+        << OpenACCClauseKind::Collapse;
 
     // Remove the value so that we don't get cascading errors in the body. The
     // caller RAII object will restore this.
     CollapseInfo.CurCollapseCount = std::nullopt;
   }
+
+  if (TileInfo.CurTileCount && *TileInfo.CurTileCount > 0) {
+    Diag(DoLoc, diag::err_acc_invalid_in_loop)
+        << /*do loop*/ 2 << OpenACCClauseKind::Tile;
+    assert(TileInfo.ActiveTile && "tile count without object?");
+    Diag(TileInfo.ActiveTile->getBeginLoc(), diag::note_acc_active_clause_here)
+        << OpenACCClauseKind::Tile;
+
+    // Remove the value so that we don't get cascading errors in the body. The
+    // caller RAII object will restore this.
+    TileInfo.CurTileCount = std::nullopt;
+  }
 }
 
 void SemaOpenACC::ActOnForStmtBegin(SourceLocation ForLoc) {
@@ -1810,7 +1860,7 @@ void SemaOpenACC::ActOnForStmtBegin(SourceLocation ForLoc) {
     return;
 
   // Enable the while/do-while checking.
-  CollapseInfo.TopLevelLoopSeen = true;
+  LoopInfo.TopLevelLoopSeen = true;
 
   if (CollapseInfo.CurCollapseCount && *CollapseInfo.CurCollapseCount > 0) {
 
@@ -1819,11 +1869,12 @@ void SemaOpenACC::ActOnForStmtBegin(SourceLocation ForLoc) {
     // or loop nest.
     // This checks for more than 1 loop at the current level, the
     // 'depth'-satisifed checking manages the 'not zero' case.
-    if (CollapseInfo.CurLevelHasLoopAlready) {
-      Diag(ForLoc, diag::err_acc_collapse_multiple_loops);
+    if (LoopInfo.CurLevelHasLoopAlready) {
+      Diag(ForLoc, diag::err_acc_clause_multiple_loops) << /*Collapse*/ 0;
       assert(CollapseInfo.ActiveCollapse && "No collapse object?");
       Diag(CollapseInfo.ActiveCollapse->getBeginLoc(),
-           diag::note_acc_collapse_clause_here);
+           diag::note_acc_active_clause_here)
+          << OpenACCClauseKind::Collapse;
     } else {
       --(*CollapseInfo.CurCollapseCount);
 
@@ -1834,13 +1885,29 @@ void SemaOpenACC::ActOnForStmtBegin(SourceLocation ForLoc) {
     }
   }
 
+  if (TileInfo.CurTileCount && *TileInfo.CurTileCount > 0) {
+    if (LoopInfo.CurLevelHasLoopAlready) {
+      Diag(ForLoc, diag::err_acc_clause_multiple_loops) << /*Tile*/ 1;
+      assert(TileInfo.ActiveTile && "No tile object?");
+      Diag(TileInfo.ActiveTile->getBeginLoc(),
+           diag::note_acc_active_clause_here)
+          << OpenACCClauseKind::Tile;
+    } else {
+      --(*TileInfo.CurTileCount);
+      // Once we've hit zero here, we know we have deep enough 'for' loops to
+      // get to the bottom.
+      if (*TileInfo.CurTileCount == 0)
+        TileInfo.TileDepthSatisfied = true;
+    }
+  }
+
   // Set this to 'false' for the body of this loop, so that the next level
   // checks independently.
-  CollapseInfo.CurLevelHasLoopAlready = false;
+  LoopInfo.CurLevelHasLoopAlready = false;
 }
 
 namespace {
-SourceLocation FindInterveningCodeInCollapseLoop(const Stmt *CurStmt) {
+SourceLocation FindInterveningCodeInLoop(const Stmt *CurStmt) {
   // We should diagnose on anything except `CompoundStmt`, `NullStmt`,
   // `ForStmt`, `CXXForRangeStmt`, since those are legal, and `WhileStmt` and
   // `DoStmt`, as those are caught as a violation elsewhere.
@@ -1858,8 +1925,7 @@ SourceLocation FindInterveningCodeInCollapseLoop(const Stmt *CurStmt) {
   // of compound statements, as long as there isn't any code inside.
   if (const auto *CS = dyn_cast<CompoundStmt>(CurStmt)) {
     for (const auto *ChildStmt : CS->children()) {
-      SourceLocation ChildStmtLoc =
-          FindInterveningCodeInCollapseLoop(ChildStmt);
+      SourceLocation ChildStmtLoc = FindInterveningCodeInLoop(ChildStmt);
       if (ChildStmtLoc.isValid())
         return ChildStmtLoc;
     }
@@ -1874,24 +1940,33 @@ void SemaOpenACC::ActOnForStmtEnd(SourceLocation ForLoc, StmtResult Body) {
   if (!getLangOpts().OpenACC)
     return;
   // Set this to 'true' so if we find another one at this level we can diagnose.
-  CollapseInfo.CurLevelHasLoopAlready = true;
+  LoopInfo.CurLevelHasLoopAlready = true;
 
   if (!Body.isUsable())
     return;
 
-  if (CollapseInfo.CurCollapseCount && *CollapseInfo.CurCollapseCount > 0 &&
-      !CollapseInfo.ActiveCollapse->hasForce()) {
-    // OpenACC 3.3: 2.9.1
-    // If the 'force' modifier does not appear, then the associated loops must
-    // be tightly nested.  If the force modifier appears, then any intervening
-    // code may be executed multiple times as needed to perform the collapse.
+  bool IsActiveCollapse = CollapseInfo.CurCollapseCount &&
+                          *CollapseInfo.CurCollapseCount > 0 &&
+                          !CollapseInfo.ActiveCollapse->hasForce();
+  bool IsActiveTile = TileInfo.CurTileCount && *TileInfo.CurTileCount > 0;
 
-    SourceLocation OtherStmtLoc = FindInterveningCodeInCollapseLoop(Body.get());
+  if (IsActiveCollapse || IsActiveTile) {
+    SourceLocation OtherStmtLoc = FindInterveningCodeInLoop(Body.get());
 
-    if (OtherStmtLoc.isValid()) {
-      Diag(OtherStmtLoc, diag::err_acc_collapse_intervening_code);
+    if (OtherStmtLoc.isValid() && IsActiveCollapse) {
+      Diag(OtherStmtLoc, diag::err_acc_intervening_code)
+          << OpenACCClauseKind::Collapse;
       Diag(CollapseInfo.ActiveCollapse->getBeginLoc(),
-           diag::note_acc_collapse_clause_here);
+           diag::note_acc_active_clause_here)
+          << OpenACCClauseKind::Collapse;
+    }
+
+    if (OtherStmtLoc.isValid() && IsActiveTile) {
+      Diag(OtherStmtLoc, diag::err_acc_intervening_code)
+          << OpenACCClauseKind::Tile;
+      Diag(TileInfo.ActiveTile->getBeginLoc(),
+           diag::note_acc_active_clause_here)
+          << OpenACCClauseKind::Tile;
     }
   }
 }
@@ -1907,11 +1982,19 @@ bool SemaOpenACC::ActOnStartStmtDirective(OpenACCDirectiveKind K,
   //
   // ALL constructs are ill-formed if there is an active 'collapse'
   if (CollapseInfo.CurCollapseCount && *CollapseInfo.CurCollapseCount > 0) {
-    Diag(StartLoc, diag::err_acc_invalid_in_collapse_loop)
-        << /*OpenACC Construct*/ 0 << K;
+    Diag(StartLoc, diag::err_acc_invalid_in_loop)
+        << /*OpenACC Construct*/ 0 << OpenACCClauseKind::Collapse << K;
     assert(CollapseInfo.ActiveCollapse && "Collapse count without object?");
     Diag(CollapseInfo.ActiveCollapse->getBeginLoc(),
-         diag::note_acc_collapse_clause_here);
+         diag::note_acc_active_clause_here)
+        << OpenACCClauseKind::Collapse;
+  }
+  if (TileInfo.CurTileCount && *TileInfo.CurTileCount > 0) {
+    Diag(StartLoc, diag::err_acc_invalid_in_loop)
+        << /*OpenACC Construct*/ 0 << OpenACCClauseKind::Tile << K;
+    assert(TileInfo.ActiveTile && "Tile count without object?");
+    Diag(TileInfo.ActiveTile->getBeginLoc(), diag::note_acc_active_clause_here)
+        << OpenACCClauseKind::Tile;
   }
 
   return diagnoseConstructAppertainment(*this, K, StartLoc, /*IsStmt=*/true);
@@ -1986,11 +2069,24 @@ StmtResult SemaOpenACC::ActOnAssociatedStmt(SourceLocation DirectiveLoc,
       return StmtError();
     }
 
-    if (!CollapseInfo.CollapseDepthSatisfied) {
-      Diag(DirectiveLoc, diag::err_acc_collapse_insufficient_loops);
-      assert(CollapseInfo.ActiveCollapse && "Collapse count without object?");
-      Diag(CollapseInfo.ActiveCollapse->getBeginLoc(),
-           diag::note_acc_collapse_clause_here);
+    if (!CollapseInfo.CollapseDepthSatisfied || !TileInfo.TileDepthSatisfied) {
+      if (!CollapseInfo.CollapseDepthSatisfied) {
+        Diag(DirectiveLoc, diag::err_acc_insufficient_loops)
+            << OpenACCClauseKind::Collapse;
+        assert(CollapseInfo.ActiveCollapse && "Collapse count without object?");
+        Diag(CollapseInfo.ActiveCollapse->getBeginLoc(),
+             diag::note_acc_active_clause_here)
+            << OpenACCClauseKind::Collapse;
+      }
+
+      if (!TileInfo.TileDepthSatisfied) {
+        Diag(DirectiveLoc, diag::err_acc_insufficient_loops)
+            << OpenACCClauseKind::Tile;
+        assert(TileInfo.ActiveTile && "Collapse count without object?");
+        Diag(TileInfo.ActiveTile->getBeginLoc(),
+             diag::note_acc_active_clause_here)
+            << OpenACCClauseKind::Tile;
+      }
       return StmtError();
     }
 

clang/test/ParserOpenACC/parse-clauses.c

@@ -1174,7 +1174,9 @@ void Tile() {
 
   // expected-error@+1{{indirection requires pointer operand ('int' invalid)}}
 #pragma acc loop tile(* returns_int() , *)
-  for(;;){}
+  for(;;){
+    for(;;);
+  }
 
 #pragma acc loop tile(*)
   for(;;){}
@@ -1184,11 +1186,19 @@ void Tile() {
 #pragma acc loop tile(5)
   for(;;){}
 #pragma acc loop tile(*, 5)
-  for(;;){}
+  for(;;){
+    for(;;);
+  }
 #pragma acc loop tile(5, *)
-  for(;;){}
+  for(;;){
+    for(;;);
+  }
 #pragma acc loop tile(5, *, 3, *)
-  for(;;){}
+  for(;;){
+    for(;;)
+      for(;;)
+        for(;;);
+  }
 }
 
 void Gang() {

clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c

@@ -166,7 +166,8 @@ void uses() {
 #pragma acc loop auto async
   for(;;);
 #pragma acc loop auto tile(1+2, 1)
-  for(;;);
+  for(;;)
+    for(;;);
   // expected-warning@+1{{OpenACC clause 'gang' not yet implemented}}
 #pragma acc loop auto gang
   for(;;);
@@ -303,7 +304,8 @@ void uses() {
 #pragma acc loop async auto
   for(;;);
 #pragma acc loop tile(1+2, 1) auto
-  for(;;);
+  for(;;)
+    for(;;);
   // expected-warning@+1{{OpenACC clause 'gang' not yet implemented}}
 #pragma acc loop gang auto
   for(;;);
@@ -441,7 +443,8 @@ void uses() {
 #pragma acc loop independent async
   for(;;);
 #pragma acc loop independent tile(1+2, 1)
-  for(;;);
+  for(;;)
+    for(;;);
   // expected-warning@+1{{OpenACC clause 'gang' not yet implemented}}
 #pragma acc loop independent gang
   for(;;);
@@ -578,7 +581,8 @@ void uses() {
 #pragma acc loop async independent
   for(;;);
 #pragma acc loop tile(1+2, 1) independent
-  for(;;);
+  for(;;)
+    for(;;);
   // expected-warning@+1{{OpenACC clause 'gang' not yet implemented}}
 #pragma acc loop gang independent
   for(;;);
@@ -725,7 +729,8 @@ void uses() {
 #pragma acc loop seq async
   for(;;);
 #pragma acc loop seq tile(1+2, 1)
-  for(;;);
+  for(;;)
+    for(;;);
   // expected-error@+1{{OpenACC 'wait' clause is not valid on 'loop' directive}}
 #pragma acc loop seq wait
   for(;;);
@@ -871,7 +876,8 @@ void uses() {
 #pragma acc loop async seq
   for(;;);
 #pragma acc loop tile(1+2, 1) seq
-  for(;;);
+  for(;;)
+    for(;;);
   // expected-error@+1{{OpenACC 'wait' clause is not valid on 'loop' directive}}
 #pragma acc loop wait seq
   for(;;);

clang/test/SemaOpenACC/loop-construct-device_type-clause.c

@@ -190,7 +190,9 @@ void uses() {
   for(;;);
 
 #pragma acc loop device_type(*) tile(*, 1)
-  for(;;);
+  for(;;)
+    for(;;);
+
   // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
 #pragma acc loop dtype(*) gang
   for(;;);

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

@@ -91,5 +91,286 @@ void negative_zero_constexpr() {
 #pragma acc loop tile(*, ConvertsOne{})
   for(;;)
     for(;;);
-
+}
+
+template<unsigned One>
+void only_for_loops_templ() {
+  // expected-note@+1{{'loop' construct is here}}
+#pragma acc loop tile(One)
+  // expected-error@+1{{OpenACC 'loop' construct can only be applied to a 'for' loop}}
+  while(true);
+
+  // expected-note@+1{{'loop' construct is here}}
+#pragma acc loop tile(One)
+  // expected-error@+1{{OpenACC 'loop' construct can only be applied to a 'for' loop}}
+  do {} while(true);
+
+  // expected-error@+1{{'tile' clause specifies a loop count greater than the number of available loops}}
+#pragma acc loop tile(One, 2) // expected-note 2{{active 'tile' clause defined here}}
+  for (;;)
+      // expected-error@+1{{while loop cannot appear in intervening code of a 'loop' with a 'tile' clause}}
+    while(true);
+
+  // expected-error@+1{{'tile' clause specifies a loop count greater than the number of available loops}}
+#pragma acc loop tile(One, 2) // expected-note 2{{active 'tile' clause defined here}}
+  for (;;)
+      // expected-error@+1{{do loop cannot appear in intervening code of a 'loop' with a 'tile' clause}}
+    do{}while(true);
+}
+
+
+void only_for_loops() {
+  // expected-note@+1{{'loop' construct is here}}
+#pragma acc loop tile(1)
+  // expected-error@+1{{OpenACC 'loop' construct can only be applied to a 'for' loop}}
+  while(true);
+
+  // expected-note@+1{{'loop' construct is here}}
+#pragma acc loop tile(1)
+  // expected-error@+1{{OpenACC 'loop' construct can only be applied to a 'for' loop}}
+  do {} while(true);
+
+  // expected-error@+1{{'tile' clause specifies a loop count greater than the number of available loops}}
+#pragma acc loop tile(1, 2) // expected-note 2{{active 'tile' clause defined here}}
+  for (;;)
+      // expected-error@+1{{while loop cannot appear in intervening code of a 'loop' with a 'tile' clause}}
+    while(true);
+
+  // expected-error@+1{{'tile' clause specifies a loop count greater than the number of available loops}}
+#pragma acc loop tile(1, 2) // expected-note 2{{active 'tile' clause defined here}}
+  for (;;)
+      // expected-error@+1{{do loop cannot appear in intervening code of a 'loop' with a 'tile' clause}}
+    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(;;);
+}
+
+template<unsigned Val>
+void depth_too_high_templ() {
+  // expected-error@+1{{'tile' clause specifies a loop count greater than the number of available loops}}
+#pragma acc loop tile (Val, *, Val) // expected-note{{active 'tile' clause defined here}}
+  for(;;)
+    for(;;);
+
+  // expected-error@+1{{'tile' clause specifies a loop count greater than the number of available loops}}
+#pragma acc loop tile (Val, *, Val) // expected-note 2{{active 'tile' clause defined here}}
+  for(;;)
+    for(;;)
+      // expected-error@+1{{while loop cannot appear in intervening code of a 'loop' with a 'tile' clause}}
+      while(true);
+
+  // expected-error@+1{{'tile' clause specifies a loop count greater than the number of available loops}}
+#pragma acc loop tile (Val, *, Val) // expected-note 2{{active 'tile' clause defined here}}
+  for(;;)
+    for(;;)
+      // expected-error@+1{{do loop cannot appear in intervening code of a 'loop' with a 'tile' clause}}
+      do{}while(true);
+
+  int Arr[Val+5];
+
+  // expected-error@+1{{'tile' clause specifies a loop count greater than the number of available loops}}
+#pragma acc loop tile (Val, *, Val) // expected-note 2{{active 'tile' clause defined here}}
+  for(;;)
+    for(auto x : Arr)
+      // expected-error@+1{{while loop cannot appear in intervening code of a 'loop' with a 'tile' clause}}
+      while(true)
+        for(;;);
+
+#pragma acc loop tile (Val, *, Val)
+  for(;;)
+    for(auto x : Arr)
+      for(;;)
+        while(true);
+}
+
+void depth_too_high() {
+  depth_too_high_templ<3>();
+
+int Arr[5];
+
+  // expected-error@+1{{'tile' clause specifies a loop count greater than the number of available loops}}
+#pragma acc loop tile (1, *, 3) // expected-note{{active 'tile' clause defined here}}
+  for(;;)
+    for(;;);
+
+  // expected-error@+1{{'tile' clause specifies a loop count greater than the number of available loops}}
+#pragma acc loop tile (1, *, 3) // expected-note 2{{active 'tile' clause defined here}}
+  for(;;)
+    for(;;)
+      // expected-error@+1{{while loop cannot appear in intervening code of a 'loop' with a 'tile' clause}}
+      while(true);
+
+  // expected-error@+1{{'tile' clause specifies a loop count greater than the number of available loops}}
+#pragma acc loop tile (1, *, 3) // expected-note 2{{active 'tile' clause defined here}}
+  for(;;)
+    for(;;)
+      // expected-error@+1{{do loop cannot appear in intervening code of a 'loop' with a 'tile' clause}}
+      do{}while(true);
+
+  // expected-error@+1{{'tile' clause specifies a loop count greater than the number of available loops}}
+#pragma acc loop tile (1, *, 3) // expected-note 2{{active 'tile' clause defined here}}
+  for(;;)
+    for(;;)
+      // expected-error@+1{{while loop cannot appear in intervening code of a 'loop' with a 'tile' clause}}
+      while(true)
+        for(;;);
+
+#pragma acc loop tile (1, *, 3)
+  for(;;)
+    for(auto x : Arr)
+      for(;;)
+        while(true);
+}
+
+template<unsigned Val>
+void not_single_loop_templ() {
+
+  int Arr[Val];
+
+#pragma acc loop tile (Val, *, 3) // expected-note{{active 'tile' clause defined here}}
+  for(;;) {
+    for (auto x : Arr)
+      for(;;);
+  // expected-error@+1{{more than one for-loop in a loop associated with OpenACC 'loop' construct with a 'tile' clause}}
+    for(;;)
+      for(;;);
+  }
+}
+
+void not_single_loop() {
+  not_single_loop_templ<3>(); // no diagnostic, was diagnosed in phase 1.
+
+  int Arr[5];
+
+#pragma acc loop tile (1, *, 3)// expected-note{{active 'tile' clause defined here}}
+  for(;;) {
+    for (auto x : Arr)
+      for(;;);
+  // expected-error@+1{{more than one for-loop in a loop associated with OpenACC 'loop' construct with a 'tile' clause}}
+    for(;;)
+      for(;;);
+  }
+}
+
+template<unsigned Val>
+void no_other_directives_templ() {
+
+  int Arr[Val];
+
+#pragma acc loop tile (Val, *, 3) // expected-note{{active 'tile' clause defined here}}
+  for(;;) {
+    for (auto x : Arr) {
+  // expected-error@+1{{OpenACC 'serial' construct cannot appear in intervening code of a 'loop' with a 'tile' clause}}
+#pragma acc serial
+      ;
+      for(;;);
+    }
+  }
+
+  // OK, in innermost
+#pragma acc loop tile (Val, *, 3)
+  for(;;) {
+    for (;;) {
+      for (auto x : Arr) {
+#pragma acc serial
+      ;
+      }
+    }
+  }
+}
+
+void no_other_directives() {
+  no_other_directives_templ<3>();
+  int Arr[5];
+
+#pragma acc loop tile (1, *, 3) // expected-note{{active 'tile' clause defined here}}
+  for(;;) {
+    for (auto x : Arr) {
+  // expected-error@+1{{OpenACC 'serial' construct cannot appear in intervening code of a 'loop' with a 'tile' clause}}
+#pragma acc serial
+      ;
+      for(;;);
+    }
+  }
+
+  // OK, in innermost
+#pragma acc loop tile (3, *, 3)
+  for(;;) {
+    for (;;) {
+      for (auto x : Arr) {
+#pragma acc serial
+      ;
+      }
+    }
+  }
+}
+
+void call();
+template<unsigned Val>
+void intervening_templ() {
+#pragma acc loop tile(1, Val, *) // expected-note{{active 'tile' clause defined here}}
+  for(;;) {
+    //expected-error@+1{{inner loops must be tightly nested inside a 'tile' clause on a 'loop' construct}}
+    call();
+    for(;;)
+      for(;;);
+  }
+
+#pragma acc loop tile(1, Val, *) // expected-note{{active 'tile' clause defined here}}
+  for(;;) {
+    //expected-error@+1{{inner loops must be tightly nested inside a 'tile' clause on a 'loop' construct}}
+    unsigned I;
+    for(;;)
+      for(;;);
+  }
+
+#pragma acc loop tile(1, Val, *)
+  for(;;) {
+    for(;;)
+      for(;;)
+        call();
+  }
+}
+
+void intervening() {
+  intervening_templ<3>();
+
+#pragma acc loop tile(1, 2, *) // expected-note{{active 'tile' clause defined here}}
+  for(;;) {
+    //expected-error@+1{{inner loops must be tightly nested inside a 'tile' clause on a 'loop' construct}}
+    call();
+    for(;;)
+      for(;;);
+  }
+
+#pragma acc loop tile(1, 2, *) // expected-note{{active 'tile' clause defined here}}
+  for(;;) {
+    //expected-error@+1{{inner loops must be tightly nested inside a 'tile' clause on a 'loop' construct}}
+    unsigned I;
+    for(;;)
+      for(;;);
+  }
+
+#pragma acc loop tile(1, 2, *)
+  for(;;) {
+    for(;;)
+      for(;;)
+        call();
+  }
+}
+
+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}}
+  // 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, 3) collapse (3)
+  for(;;) {
+    for(;;);
+  }
 }