https://github.com/erichkeane created
https://github.com/llvm/llvm-project/pull/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.
>From eb232be419fd56fe518fae9c2ddcd0606e1c4c02 Mon Sep 17 00:00:00 2001
From: erichkeane <eke...@nvidia.com>
Date: Fri, 27 Sep 2024 07:22:31 -0700
Subject: [PATCH] [OpenACC] Implement 'tile' loop count/tightly nested loop
requirement
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/Basic/DiagnosticSemaKinds.td | 22 +-
clang/include/clang/Sema/SemaOpenACC.h | 73 +++--
clang/lib/Sema/SemaOpenACC.cpp | 164 +++++++---
clang/test/ParserOpenACC/parse-clauses.c | 18 +-
...p-construct-auto_seq_independent-clauses.c | 18 +-
.../loop-construct-device_type-clause.c | 4 +-
.../loop-construct-tile-clause.cpp | 281 ++++++++++++++++++
7 files changed, 504 insertions(+), 76 deletions(-)
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td
b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 59dc81cfeb7111..583475327c5227 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/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
diff --git a/clang/include/clang/Sema/SemaOpenACC.h
b/clang/include/clang/Sema/SemaOpenACC.h
index d25c52ec3be43a..97386d2378b758 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -42,6 +42,23 @@ class SemaOpenACC : public SemaBase {
/// 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 @@ class SemaOpenACC : public SemaBase {
/// 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 @@ class SemaOpenACC : public SemaBase {
/// 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 @@ class SemaOpenACC : public SemaBase {
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 @@ class SemaOpenACC : public SemaBase {
void SetCollapseInfoBeforeAssociatedStmt(
ArrayRef<const OpenACCClause *> UnInstClauses,
ArrayRef<OpenACCClause *> Clauses);
+ void SetTileInfoBeforeAssociatedStmt(
+ ArrayRef<const OpenACCClause *> UnInstClauses,
+ ArrayRef<OpenACCClause *> Clauses);
~AssociatedStmtRAII();
};
};
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index f3840c6ffba82b..66f8029a2754b9 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/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();
}
diff --git a/clang/test/ParserOpenACC/parse-clauses.c
b/clang/test/ParserOpenACC/parse-clauses.c
index a752b2ebd18b6a..6c382379a8a7ea 100644
--- a/clang/test/ParserOpenACC/parse-clauses.c
+++ b/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() {
diff --git
a/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
b/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
index f66ce6991acb1a..3da7f0e9836be8 100644
--- a/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
+++ b/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(;;);
diff --git a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
index aa7c3201c8319c..3d77c031f42630 100644
--- a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
+++ b/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(;;);
diff --git a/clang/test/SemaOpenACC/loop-construct-tile-clause.cpp
b/clang/test/SemaOpenACC/loop-construct-tile-clause.cpp
index 28eadde7416c35..ecf8dc1109fb76 100644
--- a/clang/test/SemaOpenACC/loop-construct-tile-clause.cpp
+++ b/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(;;);
+ }
}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
