[OpenACC] 'collapse' clause 'force' tag

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -12628,6 +12628,9 @@ def err_acc_collapse_multiple_loops
 def err_acc_collapse_insufficient_loops
     : Error<"'collapse' 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 "
+            "a 'loop' construct">;
 
 // AMDGCN builtins diagnostics
 def err_amdgcn_global_load_lds_size_invalid_value : Error<"invalid size value">;

clang/lib/Sema/SemaOpenACC.cpp

@@ -1770,11 +1770,61 @@ void SemaOpenACC::ActOnForStmtBegin(SourceLocation ForLoc) {
   CollapseInfo.CurLevelHasLoopAlready = false;
 }
 
+namespace {
+SourceLocation FindInterveningCodeInCollapseLoop(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.
+  // For `CompoundStmt` we need to search inside of it.
+  if (!CurStmt ||
+      isa<ForStmt, NullStmt, ForStmt, CXXForRangeStmt, WhileStmt, DoStmt>(
+          CurStmt))
+    return SourceLocation{};
+
+  // Any other construct is an error anyway, so it has already been diagnosed.
+  if (isa<OpenACCConstructStmt>(CurStmt))
+    return SourceLocation{};
+
+  // Search inside the compound statement, this allows for arbitrary nesting
+  // 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);
+      if (ChildStmtLoc.isValid())
+        return ChildStmtLoc;
+    }
+    // Empty/not invalid compound statements are legal.
+    return SourceLocation{};
+  }
+  return CurStmt->getBeginLoc();
+}
+} // namespace
+
 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;
+
+  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.
+
+    SourceLocation OtherStmtLoc = FindInterveningCodeInCollapseLoop(Body.get());
+
+    if (OtherStmtLoc.isValid()) {
+      Diag(OtherStmtLoc, diag::err_acc_collapse_intervening_code);
+      Diag(CollapseInfo.ActiveCollapse->getBeginLoc(),
+           diag::note_acc_collapse_clause_here);
+    }
+  }
 }
 
 bool SemaOpenACC::ActOnStartStmtDirective(OpenACCDirectiveKind K,

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

@@ -336,3 +336,146 @@ void no_other_directives() {
   }
 }
 
+void call();
+
+template<unsigned Two>
+void intervening_without_force_templ() {
+  // expected-note@+1{{active 'collapse' clause defined here}}
+#pragma acc loop collapse(2)
+  for(;;) {
+    // expected-error@+1{{inner loops must be tightly nested inside a 'collapse' clause on a 'loop' construct}}
+    call();
+    for(;;){}
+  }
+
+  // expected-note@+1{{active 'collapse' clause defined here}}
+#pragma acc loop collapse(Two)
+  for(;;) {
+    // expected-error@+1{{inner loops must be tightly nested inside a 'collapse' clause on a 'loop' construct}}
+    call();
+    for(;;){}
+  }
+
+  // expected-note@+1{{active 'collapse' clause defined here}}
+#pragma acc loop collapse(2)
+  for(;;) {
+    for(;;){}
+    // expected-error@+1{{inner loops must be tightly nested inside a 'collapse' clause on a 'loop' construct}}
+    call();
+  }
+
+#pragma acc loop collapse(force:2)
+  for(;;) {
+    call();
+    for(;;){}
+  }
+
+#pragma acc loop collapse(force:Two)
+  for(;;) {
+    call();
+    for(;;){}
+  }
+
+
+#pragma acc loop collapse(force:2)
+  for(;;) {
+    for(;;){}
+    call();
+  }
+
+#pragma acc loop collapse(force:Two)
+  for(;;) {
+    for(;;){}
+    call();
+  }
+
+#pragma acc loop collapse(Two)
+  for(;;) {
+    for(;;){
+    call();
+    }
+  }
+
+#pragma acc loop collapse(Two)
+  for(;;) {
+    {
+      {
+        for(;;){
+        call();
+        }
+      }
+    }
+  }
+
+#pragma acc loop collapse(force:Two)
+  for(;;) {
+    for(;;){
+    call();
+    }
+  }
+
+  // expected-note@+1{{active 'collapse' clause defined here}}
+#pragma acc loop collapse(Two)
+  for(;;) {
+    for(;;){}
+    // expected-error@+1{{inner loops must be tightly nested inside a 'collapse' clause on a 'loop' construct}}
+    call();
+  }
+}
+
+void intervening_without_force() {
+  intervening_without_force_templ<2>(); // expected-note{{in instantiation of function template specialization}}
+  // expected-note@+1{{active 'collapse' clause defined here}}
+#pragma acc loop collapse(2)
+  for(;;) {
+    // expected-error@+1{{inner loops must be tightly nested inside a 'collapse' clause on a 'loop' construct}}
+    call();
+    for(;;){}
+  }
+
+  // expected-note@+1{{active 'collapse' clause defined here}}
+#pragma acc loop collapse(2)
+  for(;;) {
+    for(;;){}
+    // expected-error@+1{{inner loops must be tightly nested inside a 'collapse' clause on a 'loop' construct}}
+    call();
+  }
+
+  // The below two are fine, as they use the 'force' tag.
+#pragma acc loop collapse(force:2)
+  for(;;) {
+    call();
+    for(;;){}
+  }
+
+#pragma acc loop collapse(force:2)
+  for(;;) {
+    for(;;){}
+    call();
+  }
+
+#pragma acc loop collapse(2)
+  for(;;) {
+    for(;;){
+    call();
+    }
+  }
+#pragma acc loop collapse(2)
+  for(;;) {
+    {
+      {
+        for(;;){
+        call();
+        }
+      }
+    }
+  }
+
+#pragma acc loop collapse(force:2)
+  for(;;) {
+    for(;;){
+    call();
+    }
+  }
+}
+