[clang] da4b972 - [OpenACC] Enforce all 'collapse' intervening rules (#110684)
via cfe-commits
cfe-commits at lists.llvm.org
Wed Oct 2 07:10:58 PDT 2024
Author: Erich Keane
Date: 2024-10-02T07:10:53-07:00
New Revision: da4b972ef690b40344cbb366cc4bdebada260e57
URL: https://github.com/llvm/llvm-project/commit/da4b972ef690b40344cbb366cc4bdebada260e57
DIFF: https://github.com/llvm/llvm-project/commit/da4b972ef690b40344cbb366cc4bdebada260e57.diff
LOG: [OpenACC] Enforce all 'collapse' intervening rules (#110684)
'collapse' makes the 'loop' construct apply to the next N nested loops,
and 'loop' requires all associated loops be 'for' loops (or range-for).
This patch adds this restriction, plus adds a number of infrastructure
changes to permit some of the other restrictions in the future to be
implemented.
'collapse' also requires that there is not any calls to other directives
inside of the collapse region, so this also implements that limitation.
Added:
Modified:
clang/include/clang/Basic/DiagnosticSemaKinds.td
clang/include/clang/Sema/SemaOpenACC.h
clang/lib/Parse/ParseOpenACC.cpp
clang/lib/Parse/ParseStmt.cpp
clang/lib/Sema/SemaOpenACC.cpp
clang/lib/Sema/TreeTransform.h
clang/test/SemaOpenACC/loop-construct-collapse-clause.cpp
Removed:
################################################################################
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 81cd0d4ed5cf1e..880dda75ffb110 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12617,6 +12617,11 @@ def err_acc_loop_spec_conflict
def err_acc_collapse_loop_count
: Error<"OpenACC 'collapse' clause loop count must be a %select{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">;
// AMDGCN builtins diagnostics
def err_amdgcn_global_load_lds_size_invalid_value : Error<"invalid size value">;
diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h
index 839fdb79cd0ac8..43b1b1b8e870aa 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -22,6 +22,7 @@
#include "clang/Sema/Ownership.h"
#include "clang/Sema/SemaBase.h"
#include "llvm/ADT/SmallVector.h"
+#include "llvm/Support/Compiler.h"
#include <cassert>
#include <optional>
#include <utility>
@@ -41,6 +42,30 @@ class SemaOpenACC : public SemaBase {
/// above collection.
bool InsideComputeConstruct = 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
+ /// modified, and typically should be controlled by the RAII objects.
+ ///
+ /// Collapse has an 'N' count that makes it apply to a number of loops 'below'
+ /// it.
+ struct CollapseCheckingInfo {
+ 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
+ /// there is not an active collapse, or a collapse whose depth we don't know
+ /// (for example, if it is a dependent value), this should be `nullopt`,
+ /// 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;
+ } CollapseInfo{nullptr, std::nullopt, false};
+
public:
// Redeclaration of the version in OpenACCClause.h.
using DeviceTypeArgument = std::pair<IdentifierInfo *, SourceLocation>;
@@ -411,6 +436,17 @@ class SemaOpenACC : public SemaBase {
SemaOpenACC(Sema &S);
+ // Called when we encounter a 'while' statement, before looking at its 'body'.
+ void ActOnWhileStmt(SourceLocation WhileLoc);
+ // Called when we encounter a 'do' statement, before looking at its 'body'.
+ void ActOnDoStmt(SourceLocation DoLoc);
+ // Called when we encounter a 'for' statement, before looking at its 'body'.
+ void ActOnForStmtBegin(SourceLocation ForLoc);
+ // Called when we encounter a 'for' statement, after we've consumed/checked
+ // the body. This is necessary for a number of checks on the contents of the
+ // 'for' statement.
+ void ActOnForStmtEnd(SourceLocation ForLoc, StmtResult Body);
+
/// Called after parsing an OpenACC Clause so that it can be checked.
OpenACCClause *ActOnClause(ArrayRef<const OpenACCClause *> ExistingClauses,
OpenACCParsedClause &Clause);
@@ -474,6 +510,18 @@ class SemaOpenACC : public SemaBase {
/// Checks the loop depth value for a collapse clause.
ExprResult CheckCollapseLoopCount(Expr *LoopCount);
+ /// Helper type to restore the state of various 'loop' constructs when we run
+ /// into a loop (for, etc) inside the construct.
+ class LoopInConstructRAII {
+ SemaOpenACC &SemaRef;
+ CollapseCheckingInfo OldCollapseInfo;
+
+ public:
+ LoopInConstructRAII(SemaOpenACC &SemaRef)
+ : SemaRef(SemaRef), OldCollapseInfo(SemaRef.CollapseInfo) {}
+ ~LoopInConstructRAII() { SemaRef.CollapseInfo = OldCollapseInfo; }
+ };
+
/// Helper type for the registration/assignment of constructs that need to
/// 'know' about their parent constructs and hold a reference to them, such as
/// Loop needing its parent construct.
@@ -482,9 +530,15 @@ class SemaOpenACC : public SemaBase {
bool WasInsideComputeConstruct;
OpenACCDirectiveKind DirKind;
llvm::SmallVector<OpenACCLoopConstruct *> ParentlessLoopConstructs;
+ LoopInConstructRAII LoopRAII;
public:
- AssociatedStmtRAII(SemaOpenACC &, OpenACCDirectiveKind);
+ AssociatedStmtRAII(SemaOpenACC &, OpenACCDirectiveKind,
+ ArrayRef<const OpenACCClause *>,
+ ArrayRef<OpenACCClause *>);
+ void SetCollapseInfoBeforeAssociatedStmt(
+ ArrayRef<const OpenACCClause *> UnInstClauses,
+ ArrayRef<OpenACCClause *> Clauses);
~AssociatedStmtRAII();
};
};
diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp
index e66abd6873794e..87673beb34300e 100644
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@@ -1461,8 +1461,8 @@ StmtResult Parser::ParseOpenACCDirectiveStmt() {
return StmtError();
StmtResult AssocStmt;
- SemaOpenACC::AssociatedStmtRAII AssocStmtRAII(getActions().OpenACC(),
- DirInfo.DirKind);
+ SemaOpenACC::AssociatedStmtRAII AssocStmtRAII(
+ getActions().OpenACC(), DirInfo.DirKind, {}, DirInfo.Clauses);
if (doesDirectiveHaveAssociatedStmt(DirInfo.DirKind)) {
ParsingOpenACCDirectiveRAII DirScope(*this, /*Value=*/false);
ParseScope ACCScope(this, getOpenACCScopeFlags(DirInfo.DirKind));
diff --git a/clang/lib/Parse/ParseStmt.cpp b/clang/lib/Parse/ParseStmt.cpp
index 9188799fce13e6..6480e88316a7d5 100644
--- a/clang/lib/Parse/ParseStmt.cpp
+++ b/clang/lib/Parse/ParseStmt.cpp
@@ -24,6 +24,7 @@
#include "clang/Sema/Scope.h"
#include "clang/Sema/SemaCodeCompletion.h"
#include "clang/Sema/SemaObjC.h"
+#include "clang/Sema/SemaOpenACC.h"
#include "clang/Sema/SemaOpenMP.h"
#include "clang/Sema/TypoCorrection.h"
#include "llvm/ADT/STLExtras.h"
@@ -1854,6 +1855,11 @@ StmtResult Parser::ParseWhileStatement(SourceLocation *TrailingElseLoc) {
Sema::ConditionKind::Boolean, LParen, RParen))
return StmtError();
+ // OpenACC Restricts a while-loop inside of certain construct/clause
+ // combinations, so diagnose that here in OpenACC mode.
+ SemaOpenACC::LoopInConstructRAII LCR{getActions().OpenACC()};
+ getActions().OpenACC().ActOnWhileStmt(WhileLoc);
+
// C99 6.8.5p5 - In C99, the body of the while statement is a scope, even if
// there is no compound stmt. C90 does not have this clause. We only do this
// if the body isn't a compound statement to avoid push/pop in common cases.
@@ -1902,6 +1908,11 @@ StmtResult Parser::ParseDoStatement() {
ParseScope DoScope(this, ScopeFlags);
+ // OpenACC Restricts a do-while-loop inside of certain construct/clause
+ // combinations, so diagnose that here in OpenACC mode.
+ SemaOpenACC::LoopInConstructRAII LCR{getActions().OpenACC()};
+ getActions().OpenACC().ActOnDoStmt(DoLoc);
+
// C99 6.8.5p5 - In C99, the body of the do statement is a scope, even if
// there is no compound stmt. C90 does not have this clause. We only do this
// if the body isn't a compound statement to avoid push/pop in common cases.
@@ -2326,6 +2337,11 @@ StmtResult Parser::ParseForStatement(SourceLocation *TrailingElseLoc) {
}
}
+ // OpenACC Restricts a for-loop inside of certain construct/clause
+ // combinations, so diagnose that here in OpenACC mode.
+ SemaOpenACC::LoopInConstructRAII LCR{getActions().OpenACC()};
+ getActions().OpenACC().ActOnForStmtBegin(ForLoc);
+
// C99 6.8.5p5 - In C99, the body of the for statement is a scope, even if
// there is no compound stmt. C90 does not have this clause. We only do this
// if the body isn't a compound statement to avoid push/pop in common cases.
@@ -2358,6 +2374,8 @@ StmtResult Parser::ParseForStatement(SourceLocation *TrailingElseLoc) {
// Pop the body scope if needed.
InnerScope.Exit();
+ getActions().OpenACC().ActOnForStmtEnd(ForLoc, Body);
+
// Leave the for-scope.
ForScope.Exit();
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 89142b837e60a9..58ed3f3cb0777f 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -1073,19 +1073,62 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitCollapseClause(
SemaOpenACC::SemaOpenACC(Sema &S) : SemaBase(S) {}
-SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(SemaOpenACC &S,
- OpenACCDirectiveKind DK)
+SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
+ SemaOpenACC &S, OpenACCDirectiveKind DK,
+ ArrayRef<const OpenACCClause *> UnInstClauses,
+ ArrayRef<OpenACCClause *> Clauses)
: SemaRef(S), WasInsideComputeConstruct(S.InsideComputeConstruct),
- DirKind(DK) {
+ DirKind(DK), LoopRAII(SemaRef) {
// Compute constructs end up taking their 'loop'.
if (DirKind == OpenACCDirectiveKind::Parallel ||
DirKind == OpenACCDirectiveKind::Serial ||
DirKind == OpenACCDirectiveKind::Kernels) {
SemaRef.InsideComputeConstruct = true;
SemaRef.ParentlessLoopConstructs.swap(ParentlessLoopConstructs);
+ } else if (DirKind == OpenACCDirectiveKind::Loop) {
+ SetCollapseInfoBeforeAssociatedStmt(UnInstClauses, Clauses);
}
}
+void SemaOpenACC::AssociatedStmtRAII::SetCollapseInfoBeforeAssociatedStmt(
+ ArrayRef<const OpenACCClause *> UnInstClauses,
+ ArrayRef<OpenACCClause *> Clauses) {
+ // 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 =
+ 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())
+ return;
+
+ SemaRef.CollapseInfo.CurCollapseCount =
+ cast<ConstantExpr>(LoopCount)->getResultAsAPSInt();
+}
+
SemaOpenACC::AssociatedStmtRAII::~AssociatedStmtRAII() {
SemaRef.InsideComputeConstruct = WasInsideComputeConstruct;
if (DirKind == OpenACCDirectiveKind::Parallel ||
@@ -1094,6 +1137,9 @@ SemaOpenACC::AssociatedStmtRAII::~AssociatedStmtRAII() {
assert(SemaRef.ParentlessLoopConstructs.empty() &&
"Didn't consume loop construct list?");
SemaRef.ParentlessLoopConstructs.swap(ParentlessLoopConstructs);
+ } else if (DirKind == OpenACCDirectiveKind::Loop) {
+ // Nothing really to do here, the LoopInConstruct should handle restorations
+ // correctly.
}
}
@@ -1646,10 +1692,78 @@ ExprResult SemaOpenACC::CheckCollapseLoopCount(Expr *LoopCount) {
ConstantExpr::Create(getASTContext(), LoopCount, APValue{*ICE})};
}
+void SemaOpenACC::ActOnWhileStmt(SourceLocation WhileLoc) {
+ if (!getLangOpts().OpenACC)
+ return;
+
+ if (!CollapseInfo.TopLevelLoopSeen)
+ return;
+
+ if (CollapseInfo.CurCollapseCount && *CollapseInfo.CurCollapseCount > 0) {
+ Diag(WhileLoc, diag::err_acc_invalid_in_collapse_loop) << /*while loop*/ 1;
+ assert(CollapseInfo.ActiveCollapse && "Collapse count without object?");
+ Diag(CollapseInfo.ActiveCollapse->getBeginLoc(),
+ diag::note_acc_collapse_clause_here);
+
+ // 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;
+ }
+}
+
+void SemaOpenACC::ActOnDoStmt(SourceLocation DoLoc) {
+ if (!getLangOpts().OpenACC)
+ return;
+
+ if (!CollapseInfo.TopLevelLoopSeen)
+ return;
+
+ if (CollapseInfo.CurCollapseCount && *CollapseInfo.CurCollapseCount > 0) {
+ Diag(DoLoc, diag::err_acc_invalid_in_collapse_loop) << /*do loop*/ 2;
+ assert(CollapseInfo.ActiveCollapse && "Collapse count without object?");
+ Diag(CollapseInfo.ActiveCollapse->getBeginLoc(),
+ diag::note_acc_collapse_clause_here);
+
+ // 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;
+ }
+}
+
+void SemaOpenACC::ActOnForStmtBegin(SourceLocation ForLoc) {
+ if (!getLangOpts().OpenACC)
+ return;
+
+ // Enable the while/do-while checking.
+ CollapseInfo.TopLevelLoopSeen = true;
+
+ if (CollapseInfo.CurCollapseCount && *CollapseInfo.CurCollapseCount > 0)
+ --(*CollapseInfo.CurCollapseCount);
+}
+
+void SemaOpenACC::ActOnForStmtEnd(SourceLocation ForLoc, StmtResult Body) {
+ if (!getLangOpts().OpenACC)
+ return;
+}
+
bool SemaOpenACC::ActOnStartStmtDirective(OpenACCDirectiveKind K,
SourceLocation StartLoc) {
SemaRef.DiscardCleanupsInEvaluationContext();
SemaRef.PopExpressionEvaluationContext();
+
+ // OpenACC 3.3 2.9.1:
+ // Intervening code must not contain other OpenACC directives or calls to API
+ // routines.
+ //
+ // 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;
+ assert(CollapseInfo.ActiveCollapse && "Collapse count without object?");
+ Diag(CollapseInfo.ActiveCollapse->getBeginLoc(),
+ diag::note_acc_collapse_clause_here);
+ }
+
return diagnoseConstructAppertainment(*this, K, StartLoc, /*IsStmt=*/true);
}
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index 6fdb18d51acef9..04308b2b57d8a0 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -8200,6 +8200,11 @@ TreeTransform<Derived>::TransformWhileStmt(WhileStmt *S) {
if (Cond.isInvalid())
return StmtError();
+ // OpenACC Restricts a while-loop inside of certain construct/clause
+ // combinations, so diagnose that here in OpenACC mode.
+ SemaOpenACC::LoopInConstructRAII LCR{SemaRef.OpenACC()};
+ SemaRef.OpenACC().ActOnWhileStmt(S->getBeginLoc());
+
// Transform the body
StmtResult Body = getDerived().TransformStmt(S->getBody());
if (Body.isInvalid())
@@ -8217,6 +8222,11 @@ TreeTransform<Derived>::TransformWhileStmt(WhileStmt *S) {
template<typename Derived>
StmtResult
TreeTransform<Derived>::TransformDoStmt(DoStmt *S) {
+ // OpenACC Restricts a do-loop inside of certain construct/clause
+ // combinations, so diagnose that here in OpenACC mode.
+ SemaOpenACC::LoopInConstructRAII LCR{SemaRef.OpenACC()};
+ SemaRef.OpenACC().ActOnDoStmt(S->getBeginLoc());
+
// Transform the body
StmtResult Body = getDerived().TransformStmt(S->getBody());
if (Body.isInvalid())
@@ -8270,11 +8280,18 @@ TreeTransform<Derived>::TransformForStmt(ForStmt *S) {
if (S->getInc() && !FullInc.get())
return StmtError();
+ // OpenACC Restricts a for-loop inside of certain construct/clause
+ // combinations, so diagnose that here in OpenACC mode.
+ SemaOpenACC::LoopInConstructRAII LCR{SemaRef.OpenACC()};
+ SemaRef.OpenACC().ActOnForStmtBegin(S->getBeginLoc());
+
// Transform the body
StmtResult Body = getDerived().TransformStmt(S->getBody());
if (Body.isInvalid())
return StmtError();
+ SemaRef.OpenACC().ActOnForStmtEnd(S->getBeginLoc(), Body);
+
if (!getDerived().AlwaysRebuild() &&
Init.get() == S->getInit() &&
Cond.get() == std::make_pair(S->getConditionVariable(), S->getCond()) &&
@@ -9008,10 +9025,17 @@ TreeTransform<Derived>::TransformCXXForRangeStmt(CXXForRangeStmt *S) {
}
}
+ // OpenACC Restricts a while-loop inside of certain construct/clause
+ // combinations, so diagnose that here in OpenACC mode.
+ SemaOpenACC::LoopInConstructRAII LCR{SemaRef.OpenACC()};
+ SemaRef.OpenACC().ActOnForStmtBegin(S->getBeginLoc());
+
StmtResult Body = getDerived().TransformStmt(S->getBody());
if (Body.isInvalid())
return StmtError();
+ SemaRef.OpenACC().ActOnForStmtEnd(S->getBeginLoc(), Body);
+
// Body has changed but we didn't rebuild the for-range statement. Rebuild
// it now so we have a new statement to attach the body to.
if (Body.get() != S->getBody() && NewStmt.get() == S) {
@@ -11894,8 +11918,9 @@ StmtResult TreeTransform<Derived>::TransformOpenACCComputeConstruct(
return StmtError();
// Transform Structured Block.
- SemaOpenACC::AssociatedStmtRAII AssocStmtRAII(getSema().OpenACC(),
- C->getDirectiveKind());
+ SemaOpenACC::AssociatedStmtRAII AssocStmtRAII(
+ getSema().OpenACC(), C->getDirectiveKind(), C->clauses(),
+ TransformedClauses);
StmtResult StrBlock = getDerived().TransformStmt(C->getStructuredBlock());
StrBlock = getSema().OpenACC().ActOnAssociatedStmt(
C->getBeginLoc(), C->getDirectiveKind(), StrBlock);
@@ -11920,8 +11945,9 @@ TreeTransform<Derived>::TransformOpenACCLoopConstruct(OpenACCLoopConstruct *C) {
return StmtError();
// Transform Loop.
- SemaOpenACC::AssociatedStmtRAII AssocStmtRAII(getSema().OpenACC(),
- C->getDirectiveKind());
+ SemaOpenACC::AssociatedStmtRAII AssocStmtRAII(
+ getSema().OpenACC(), C->getDirectiveKind(), C->clauses(),
+ TransformedClauses);
StmtResult Loop = getDerived().TransformStmt(C->getLoop());
Loop = getSema().OpenACC().ActOnAssociatedStmt(C->getBeginLoc(),
C->getDirectiveKind(), Loop);
diff --git a/clang/test/SemaOpenACC/loop-construct-collapse-clause.cpp b/clang/test/SemaOpenACC/loop-construct-collapse-clause.cpp
index 9c1e577773e8f8..6e4a7e9468aef8 100644
--- a/clang/test/SemaOpenACC/loop-construct-collapse-clause.cpp
+++ b/clang/test/SemaOpenACC/loop-construct-collapse-clause.cpp
@@ -115,3 +115,160 @@ void negative_constexpr(int i) {
negative_constexpr_templ<int, 1>(); // #NCET1
}
+template<typename T, unsigned Three>
+void not_single_loop_templ() {
+ T Arr[5];
+ // expected-note at +1{{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 at +1{{active 'collapse' clause defined here}}
+#pragma acc loop collapse(Three)
+ for(;;) {
+ for(;;){
+ do{}while(true); // expected-error{{do loop cannot appear in intervening code of a 'loop' with a 'collapse' clause}}
+ }
+ }
+
+#pragma acc loop collapse(Three)
+ for(;;) {
+ for(;;){
+ for(;;){
+ do{}while(true);
+ }
+ }
+ }
+ // expected-note at +1{{active 'collapse' clause defined here}}
+#pragma acc loop collapse(Three)
+ 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}}
+ }
+ }
+
+#pragma acc loop collapse(Three)
+ for(auto x : Arr) {
+ for(auto y: Arr) {
+ for(auto z: Arr) {
+ do{}while(true);
+ }
+ }
+ }
+}
+
+void not_single_loop() {
+ not_single_loop_templ<int, 3>(); // expected-note{{in instantiation of function template}}
+
+ // expected-note at +1{{active 'collapse' clause defined here}}
+#pragma acc loop collapse(3)
+ for(;;) {
+ for(;;){
+ for(;;);
+ }
+ while(true); // expected-error{{while loop cannot appear in intervening code of a 'loop' with a 'collapse' clause}}
+ }
+
+ // expected-note at +1{{active 'collapse' clause defined here}}
+#pragma acc loop collapse(3)
+ for(;;) {
+ for(;;){
+ for(;;);
+ }
+ do{}while(true); // expected-error{{do loop cannot appear in intervening code of a 'loop' with a 'collapse' clause}}
+ }
+
+ // expected-note at +1{{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 at +1{{active 'collapse' clause defined here}}
+#pragma acc loop collapse(3)
+ for(;;) {
+ for(;;){
+ do{}while(true); // expected-error{{do loop cannot appear in intervening code of a 'loop' with a 'collapse' clause}}
+ }
+ }
+
+#pragma acc loop collapse(2)
+ for(;;) {
+ for(;;){
+ do{}while(true);
+ }
+ }
+#pragma acc loop collapse(2)
+ for(;;) {
+ for(;;){
+ while(true);
+ }
+ }
+
+ int Arr[5];
+
+ // expected-note at +1{{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}}
+ }
+ }
+
+}
+
+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}}
+ for(;;) {
+ for(;;) {
+ // expected-error at +1{{OpenACC 'serial' construct cannot appear in intervening code of a 'loop' with a 'collapse' clause}}
+#pragma acc serial
+ ;
+ }
+ }
+ }
+ }
+#pragma acc loop collapse(Two)// expected-note{{active 'collapse' clause defined here}}
+ for(;;) {
+ for (;;) { // last loop associated with the top level.
+#pragma acc loop collapse(Three)
+ for(;;) {
+ for(;;) {
+ }
+ }
+ }
+ // expected-error at +1{{OpenACC 'serial' construct cannot appear in intervening code of a 'loop' with a 'collapse' clause}}
+#pragma acc serial
+ ;
+ }
+}
+
+void no_other_directives() {
+ no_other_directives<2,3>(); // expected-note{{in instantiation of function template specialization}}
+
+ // Ok, not inside the intervening list
+#pragma acc loop collapse(2)
+ for(;;) {
+ for(;;) {
+#pragma acc data // expected-warning{{OpenACC construct 'data' not yet implemented}}
+ }
+ }
+ // expected-note at +1{{active 'collapse' clause defined here}}
+#pragma acc loop collapse(2)
+ for(;;) {
+ // expected-error at +1{{OpenACC 'data' construct cannot appear in intervening code of a 'loop' with a 'collapse' clause}}
+#pragma acc data // expected-warning{{OpenACC construct 'data' not yet implemented}}
+ for(;;) {
+ }
+ }
+}
+
More information about the cfe-commits
mailing list