[OpenACC] Check Loop counts for 'collapse' clause.

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -12622,6 +12622,12 @@ def err_acc_invalid_in_collapse_loop
             "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
+    : 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 "
+            "of available loops">;
 
 // AMDGCN builtins diagnostics
 def err_amdgcn_global_load_lds_size_invalid_value : Error<"invalid size value">;

clang/include/clang/Sema/SemaOpenACC.h

@@ -64,7 +64,20 @@ class SemaOpenACC : public SemaBase {
     /// 'collapse inner loop not a 'for' loop' diagnostic.
     LLVM_PREFERRED_TYPE(bool)
     unsigned TopLevelLoopSeen : 1;
-  } CollapseInfo{nullptr, std::nullopt, false};
+
+    /// 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};
 
 public:
   // Redeclaration of the version in OpenACCClause.h.
@@ -515,11 +528,26 @@ class SemaOpenACC : public SemaBase {
   class LoopInConstructRAII {
     SemaOpenACC &SemaRef;
     CollapseCheckingInfo OldCollapseInfo;
+    bool PreserveDepth;
 
   public:
-    LoopInConstructRAII(SemaOpenACC &SemaRef)
-        : SemaRef(SemaRef), OldCollapseInfo(SemaRef.CollapseInfo) {}
-    ~LoopInConstructRAII() { SemaRef.CollapseInfo = OldCollapseInfo; }
+    LoopInConstructRAII(SemaOpenACC &SemaRef, bool PreserveDepth = true)
+        : SemaRef(SemaRef), OldCollapseInfo(SemaRef.CollapseInfo),
+          PreserveDepth(PreserveDepth) {}
+    ~LoopInConstructRAII() {
+      // The associated-statement level of this should NOT preserve this, as it
+      // is a new construct, but other loop uses need to preserve the depth so
+      // it makes it to the 'top level' for diagnostics.
+      bool CollapseDepthSatisified =
+          PreserveDepth ? SemaRef.CollapseInfo.CollapseDepthSatisfied
+                        : OldCollapseInfo.CollapseDepthSatisfied;
+      bool CurLevelHasLoopAlready =
+          PreserveDepth ? SemaRef.CollapseInfo.CurLevelHasLoopAlready
+                        : OldCollapseInfo.CurLevelHasLoopAlready;
+      SemaRef.CollapseInfo = OldCollapseInfo;
+      SemaRef.CollapseInfo.CollapseDepthSatisfied = CollapseDepthSatisified;
+      SemaRef.CollapseInfo.CurLevelHasLoopAlready = CurLevelHasLoopAlready;
+    }
   };
 
   /// Helper type for the registration/assignment of constructs that need to

clang/lib/Sema/SemaOpenACC.cpp

@@ -1078,7 +1078,7 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
     ArrayRef<const OpenACCClause *> UnInstClauses,
     ArrayRef<OpenACCClause *> Clauses)
     : SemaRef(S), WasInsideComputeConstruct(S.InsideComputeConstruct),
-      DirKind(DK), LoopRAII(SemaRef) {
+      DirKind(DK), LoopRAII(SemaRef, /*PreserveDepth=*/false) {
   // Compute constructs end up taking their 'loop'.
   if (DirKind == OpenACCDirectiveKind::Parallel ||
       DirKind == OpenACCDirectiveKind::Serial ||
@@ -1093,6 +1093,11 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
 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.CollapseInfo.CurLevelHasLoopAlready = false;
+  SemaRef.CollapseInfo.CollapseDepthSatisfied = 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
@@ -1125,6 +1130,7 @@ void SemaOpenACC::AssociatedStmtRAII::SetCollapseInfoBeforeAssociatedStmt(
            ->isInstantiationDependent())
     return;
 
+  SemaRef.CollapseInfo.CollapseDepthSatisfied = false;
   SemaRef.CollapseInfo.CurCollapseCount =
       cast<ConstantExpr>(LoopCount)->getResultAsAPSInt();
 }
@@ -1737,13 +1743,38 @@ void SemaOpenACC::ActOnForStmtBegin(SourceLocation ForLoc) {
   // Enable the while/do-while checking.
   CollapseInfo.TopLevelLoopSeen = true;
 
-  if (CollapseInfo.CurCollapseCount && *CollapseInfo.CurCollapseCount > 0)
-    --(*CollapseInfo.CurCollapseCount);
+  if (CollapseInfo.CurCollapseCount && *CollapseInfo.CurCollapseCount > 0) {
+
+    // OpenACC 3.3 2.9.1:
+    // Each associated loop, except the innermost, must contain exactly one loop
+    // 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);
+      assert(CollapseInfo.ActiveCollapse && "No collapse object?");
+      Diag(CollapseInfo.ActiveCollapse->getBeginLoc(),
+           diag::note_acc_collapse_clause_here);
+    } else {
+      --(*CollapseInfo.CurCollapseCount);
+
+      // Once we've hit zero here, we know we have deep enough 'for' loops to
+      // get to the bottom.
+      if (*CollapseInfo.CurCollapseCount == 0)
+        CollapseInfo.CollapseDepthSatisfied = true;
+    }
+  }
+
+  // Set this to 'false' for the body of this loop, so that the next level
+  // checks independently.
+  CollapseInfo.CurLevelHasLoopAlready = false;
 }
 
 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;
 }
 
 bool SemaOpenACC::ActOnStartStmtDirective(OpenACCDirectiveKind K,
@@ -1827,12 +1858,23 @@ StmtResult SemaOpenACC::ActOnAssociatedStmt(SourceLocation DirectiveLoc,
     // the 'structured block'.
     return AssocStmt;
   case OpenACCDirectiveKind::Loop:
-    if (AssocStmt.isUsable() &&
-        !isa<CXXForRangeStmt, ForStmt>(AssocStmt.get())) {
+    if (!AssocStmt.isUsable())
+      return StmtError();
+
+    if (!isa<CXXForRangeStmt, ForStmt>(AssocStmt.get())) {
       Diag(AssocStmt.get()->getBeginLoc(), diag::err_acc_loop_not_for_loop);
       Diag(DirectiveLoc, diag::note_acc_construct_here) << K;
       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);
+      return StmtError();
+    }
+
     // TODO OpenACC: 2.9 ~ line 2010 specifies that the associated loop has some
     // restrictions when there is a 'seq' clause in place. We probably need to
     // implement that, including piping in the clauses here.

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

@@ -115,18 +115,51 @@ void negative_constexpr(int i) {
   negative_constexpr_templ<int, 1>(); // #NCET1
 }
 
+template<unsigned Val>
+void depth_too_high_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(Val)
+  for(;;)
+    for(;;);
+}
+
+void depth_too_high() {
+  depth_too_high_templ<3>(); // expected-note{{in instantiation of function template specialization}}
+
+  // 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)
+  for(;;)
+    for(;;);
+
+  // 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(three())
+  for(;;)
+    for(;;);
+
+  // 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(ConvertsThree{})
+  for(;;)
+    for(;;);
+}
+
 template<typename T, unsigned Three>
 void not_single_loop_templ() {
   T Arr[5];
-  // expected-note@+1{{active 'collapse' clause defined here}}
+  // expected-error@+2{{'collapse' clause specifies a loop count greater than the number of available loops}}
+  // expected-note@+1 2{{active 'collapse' clause defined here}}
 #pragma acc loop collapse(3)
   for(auto x : Arr) {
     for(auto y : Arr){
       do{}while(true); // expected-error{{do loop cannot appear in intervening code of a 'loop' with a 'collapse' clause}}
     }
   }
 
-  // expected-note@+1{{active 'collapse' clause defined here}}
+  // expected-error@+2{{'collapse' clause specifies a loop count greater than the number of available loops}}
+  // expected-note@+1 2{{active 'collapse' clause defined here}}
 #pragma acc loop collapse(Three)
   for(;;) {
     for(;;){
@@ -142,7 +175,8 @@ void not_single_loop_templ() {
       }
     }
   }
-  // expected-note@+1{{active 'collapse' clause defined here}}
+  // expected-error@+2{{'collapse' clause specifies a loop count greater than the number of available loops}}
+  // expected-note@+1 2{{active 'collapse' clause defined here}}
 #pragma acc loop collapse(Three)
   for(auto x : Arr) {
     for(auto y: Arr) {
@@ -181,15 +215,16 @@ void not_single_loop() {
     do{}while(true); // expected-error{{do loop cannot appear in intervening code of a 'loop' with a 'collapse' clause}}
   }
 
-  // expected-note@+1{{active 'collapse' clause defined here}}
+  // expected-error@+2{{'collapse' clause specifies a loop count greater than the number of available loops}}
+  // expected-note@+1 2{{active 'collapse' clause defined here}}
 #pragma acc loop collapse(3)
   for(;;) {
     for(;;){
       while(true); // expected-error{{while loop cannot appear in intervening code of a 'loop' with a 'collapse' clause}}
     }
   }
-
-  // expected-note@+1{{active 'collapse' clause defined here}}
+  // expected-error@+2{{'collapse' clause specifies a loop count greater than the number of available loops}}
+  // expected-note@+1 2{{active 'collapse' clause defined here}}
 #pragma acc loop collapse(3)
   for(;;) {
     for(;;){
@@ -211,23 +246,51 @@ void not_single_loop() {
   }
 
   int Arr[5];
-
-  // expected-note@+1{{active 'collapse' clause defined here}}
+  // expected-error@+2{{'collapse' clause specifies a loop count greater than the number of available loops}}
+  // expected-note@+1 2{{active 'collapse' clause defined here}}
 #pragma acc loop collapse(3)
   for(auto x : Arr) {
     for(auto y : Arr){
       do{}while(true); // expected-error{{do loop cannot appear in intervening code of a 'loop' with a 'collapse' clause}}
     }
   }
 
+  // expected-note@+1 {{active 'collapse' clause defined here}}
+#pragma acc loop collapse(3)
+  for (;;) {
+    for (;;) {
+      for(;;);
+    }
+    // expected-error@+1{{more than one for-loop in a loop associated with OpenACC 'loop' construct with a 'collapse' clause}}
+    for(;;);
+  }
+
+  // expected-note@+1 {{active 'collapse' clause defined here}}
+#pragma acc loop collapse(3)
+  for (;;) {
+    for (;;) {
+      for(;;);
+    // expected-error@+1{{more than one for-loop in a loop associated with OpenACC 'loop' construct with a 'collapse' clause}}
+      for(;;);
+    }
+  }
+
+  for(;;);
+#pragma acc loop collapse(3)
+  for (;;) {
+    for (;;) {
+      for (;;);
+    }
+  }
 }
 
 template<unsigned Two, unsigned Three>
 void no_other_directives() {
 #pragma acc loop collapse(Two)
   for(;;) {
     for (;;) { // last loop associated with the top level.
-#pragma acc loop collapse(Three) // expected-note{{active 'collapse' clause defined here}}
+    // expected-error@+1{{'collapse' clause specifies a loop count greater than the number of available loops}}
+#pragma acc loop collapse(Three) // expected-note 2{{active 'collapse' clause defined here}}
       for(;;) {
         for(;;) {
     // expected-error@+1{{OpenACC 'serial' construct cannot appear in intervening code of a 'loop' with a 'collapse' clause}}
@@ -243,6 +306,7 @@ void no_other_directives() {
 #pragma acc loop collapse(Three)
       for(;;) {
         for(;;) {
+          for(;;);
         }
       }
     }