https://github.com/erichkeane created
https://github.com/llvm/llvm-project/pull/112006
The 'gang' clause is used to specify parallel execution of loops, thus has some
complicated rules depending on the 'loop's associated compute construct. This
patch implements all of those.
>From 8d6dd131cc50e747fba0c7c8a67d2bb8a4f2f231 Mon Sep 17 00:00:00 2001
From: erichkeane <eke...@nvidia.com>
Date: Tue, 8 Oct 2024 12:28:29 -0700
Subject: [PATCH] [OpenACC] Implement loop 'gang' clause.
The 'gang' clause is used to specify parallel execution of loops, thus
has some complicated rules depending on the 'loop's associated compute
construct. This patch implements all of those.
---
clang/include/clang/AST/OpenACCClause.h | 60 ++--
.../clang/Basic/DiagnosticSemaKinds.td | 19 +
clang/include/clang/Basic/OpenACCClauses.def | 1 +
clang/include/clang/Basic/OpenACCKinds.h | 29 ++
clang/include/clang/Parse/Parser.h | 12 +-
clang/include/clang/Sema/SemaOpenACC.h | 80 ++++-
clang/lib/AST/OpenACCClause.cpp | 50 ++-
clang/lib/AST/StmtProfile.cpp | 6 +
clang/lib/AST/TextNodeDumper.cpp | 11 +
clang/lib/Parse/ParseOpenACC.cpp | 50 ++-
clang/lib/Sema/SemaOpenACC.cpp | 276 +++++++++++++--
clang/lib/Sema/TreeTransform.h | 23 ++
clang/lib/Serialization/ASTReader.cpp | 13 +-
clang/lib/Serialization/ASTWriter.cpp | 11 +-
.../AST/ast-print-openacc-loop-construct.cpp | 82 +++++
clang/test/ParserOpenACC/parse-clauses.c | 50 ++-
.../compute-construct-device_type-clause.c | 3 +-
...p-construct-auto_seq_independent-clauses.c | 15 +-
.../loop-construct-device_type-clause.c | 1 -
.../SemaOpenACC/loop-construct-gang-ast.cpp | 330 +++++++++++++++++
.../loop-construct-gang-clause.cpp | 335 ++++++++++++++++++
clang/tools/libclang/CIndex.cpp | 5 +
22 files changed, 1336 insertions(+), 126 deletions(-)
create mode 100644 clang/test/SemaOpenACC/loop-construct-gang-ast.cpp
create mode 100644 clang/test/SemaOpenACC/loop-construct-gang-clause.cpp
diff --git a/clang/include/clang/AST/OpenACCClause.h
b/clang/include/clang/AST/OpenACCClause.h
index b500acc768e55a..f3a09eb651458d 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -119,32 +119,6 @@ class OpenACCSeqClause : public OpenACCClause {
}
};
-// Not yet implemented, but the type name is necessary for 'seq' diagnostics,
so
-// this provides a basic, do-nothing implementation. We still need to add this
-// type to the visitors/etc, as well as get it to take its proper arguments.
-class OpenACCGangClause : public OpenACCClause {
-protected:
- OpenACCGangClause(SourceLocation BeginLoc, SourceLocation EndLoc)
- : OpenACCClause(OpenACCClauseKind::Gang, BeginLoc, EndLoc) {
- llvm_unreachable("Not yet implemented");
- }
-
-public:
- static bool classof(const OpenACCClause *C) {
- return C->getClauseKind() == OpenACCClauseKind::Gang;
- }
-
- static OpenACCGangClause *
- Create(const ASTContext &Ctx, SourceLocation BeginLoc, SourceLocation
EndLoc);
-
- child_range children() {
- return child_range(child_iterator(), child_iterator());
- }
- const_child_range children() const {
- return const_child_range(const_child_iterator(), const_child_iterator());
- }
-};
-
// Not yet implemented, but the type name is necessary for 'seq' diagnostics,
so
// this provides a basic, do-nothing implementation. We still need to add this
// type to the visitors/etc, as well as get it to take its proper arguments.
@@ -177,7 +151,7 @@ class OpenACCVectorClause : public OpenACCClause {
class OpenACCWorkerClause : public OpenACCClause {
protected:
OpenACCWorkerClause(SourceLocation BeginLoc, SourceLocation EndLoc)
- : OpenACCClause(OpenACCClauseKind::Gang, BeginLoc, EndLoc) {
+ : OpenACCClause(OpenACCClauseKind::Worker, BeginLoc, EndLoc) {
llvm_unreachable("Not yet implemented");
}
@@ -535,6 +509,38 @@ class OpenACCClauseWithSingleIntExpr : public
OpenACCClauseWithExprs {
Expr *getIntExpr() { return hasIntExpr() ? getExprs()[0] : nullptr; };
};
+class OpenACCGangClause final
+ : public OpenACCClauseWithExprs,
+ public llvm::TrailingObjects<OpenACCGangClause, Expr *, OpenACCGangKind>
{
+protected:
+ OpenACCGangClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
+ ArrayRef<OpenACCGangKind> GangKinds,
+ ArrayRef<Expr *> IntExprs, SourceLocation EndLoc);
+
+ OpenACCGangKind getGangKind(unsigned I) const {
+ return getTrailingObjects<OpenACCGangKind>()[I];
+ }
+
+public:
+ static bool classof(const OpenACCClause *C) {
+ return C->getClauseKind() == OpenACCClauseKind::Gang;
+ }
+
+ size_t numTrailingObjects(OverloadToken<Expr *>) const {
+ return getNumExprs();
+ }
+
+ unsigned getNumExprs() const { return getExprs().size(); }
+ std::pair<OpenACCGangKind, const Expr *> getExpr(unsigned I) const {
+ return {getGangKind(I), getExprs()[I]};
+ }
+
+ static OpenACCGangClause *
+ Create(const ASTContext &Ctx, SourceLocation BeginLoc,
+ SourceLocation LParenLoc, ArrayRef<OpenACCGangKind> GangKinds,
+ ArrayRef<Expr *> IntExprs, SourceLocation EndLoc);
+};
+
class OpenACCNumWorkersClause : public OpenACCClauseWithSingleIntExpr {
OpenACCNumWorkersClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
Expr *IntExpr, SourceLocation EndLoc);
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td
b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 41cdd09e971651..3c62a017005e59 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12576,6 +12576,7 @@ def err_acc_duplicate_clause_disallowed
: Error<"OpenACC '%1' clause cannot appear more than once on a '%0' "
"directive">;
def note_acc_previous_clause_here : Note<"previous clause is here">;
+def note_acc_previous_expr_here : Note<"previous expression is here">;
def err_acc_branch_in_out_compute_construct
: Error<"invalid %select{branch|return|throw}0 %select{out of|into}1 "
"OpenACC Compute Construct">;
@@ -12682,6 +12683,24 @@ def err_acc_insufficient_loops
def err_acc_intervening_code
: Error<"inner loops must be tightly nested inside a '%0' clause on "
"a 'loop' construct">;
+def err_acc_gang_multiple_elt
+ : Error<"OpenACC 'gang' clause may have at most one %select{unnamed or "
+ "'num'|'dim'|'static'}0 argument">;
+def err_acc_gang_arg_invalid
+ : Error<"'%0' argument on 'gang' clause is not permitted on a%select{n "
+ "orphaned|||}1 'loop' construct %select{|associated with a "
+ "'parallel' compute construct|associated with a 'kernels' compute "
+ "construct|associated with a 'serial' compute construct}1">;
+def err_acc_gang_dim_value
+ : Error<"argument to 'gang' clause dimension must be %select{a constant "
+ "expression|1, 2, or 3: evaluated to %1}0">;
+def err_acc_gang_num_gangs_conflict
+ : Error<"'num' argument to 'gang' clause not allowed on a 'loop' construct
"
+ "associated with a 'kernels' construct that has a 'num_gangs' "
+ "clause">;
+def err_acc_gang_inside_gang
+ : Error<"loop with a 'gang' clause may not exist in the region of a 'gang'
"
+ "clause on a 'kernels' compute construct">;
// AMDGCN builtins diagnostics
def err_amdgcn_global_load_lds_size_invalid_value : Error<"invalid size
value">;
diff --git a/clang/include/clang/Basic/OpenACCClauses.def
b/clang/include/clang/Basic/OpenACCClauses.def
index a380e5ae69c418..2a098de31eb618 100644
--- a/clang/include/clang/Basic/OpenACCClauses.def
+++ b/clang/include/clang/Basic/OpenACCClauses.def
@@ -42,6 +42,7 @@ VISIT_CLAUSE(DevicePtr)
VISIT_CLAUSE(DeviceType)
CLAUSE_ALIAS(DType, DeviceType, false)
VISIT_CLAUSE(FirstPrivate)
+VISIT_CLAUSE(Gang)
VISIT_CLAUSE(If)
VISIT_CLAUSE(Independent)
VISIT_CLAUSE(NoCreate)
diff --git a/clang/include/clang/Basic/OpenACCKinds.h
b/clang/include/clang/Basic/OpenACCKinds.h
index c4dfe3bedc13a7..3f48ebca708a42 100644
--- a/clang/include/clang/Basic/OpenACCKinds.h
+++ b/clang/include/clang/Basic/OpenACCKinds.h
@@ -550,6 +550,35 @@ inline llvm::raw_ostream &operator<<(llvm::raw_ostream
&Out,
OpenACCReductionOperator Op) {
return printOpenACCReductionOperator(Out, Op);
}
+
+enum class OpenACCGangKind : uint8_t {
+ /// num:
+ Num,
+ /// dim:
+ Dim,
+ /// static:
+ Static
+};
+
+template <typename StreamTy>
+inline StreamTy &printOpenACCGangKind(StreamTy &Out, OpenACCGangKind GK) {
+ switch (GK) {
+ case OpenACCGangKind::Num:
+ return Out << "num";
+ case OpenACCGangKind::Dim:
+ return Out << "dim";
+ case OpenACCGangKind::Static:
+ return Out << "static";
+ }
+}
+inline const StreamingDiagnostic &operator<<(const StreamingDiagnostic &Out,
+ OpenACCGangKind Op) {
+ return printOpenACCGangKind(Out, Op);
+}
+inline llvm::raw_ostream &operator<<(llvm::raw_ostream &Out,
+ OpenACCGangKind Op) {
+ return printOpenACCGangKind(Out, Op);
+}
} // namespace clang
#endif // LLVM_CLANG_BASIC_OPENACCKINDS_H
diff --git a/clang/include/clang/Parse/Parser.h
b/clang/include/clang/Parse/Parser.h
index dbcb545058a026..045ee754a242b3 100644
--- a/clang/include/clang/Parse/Parser.h
+++ b/clang/include/clang/Parse/Parser.h
@@ -3797,9 +3797,15 @@ class Parser : public CodeCompletionHandler {
bool ParseOpenACCSizeExprList(OpenACCClauseKind CK,
llvm::SmallVectorImpl<Expr *> &SizeExprs);
/// Parses a 'gang-arg-list', used for the 'gang' clause.
- bool ParseOpenACCGangArgList(SourceLocation GangLoc);
- /// Parses a 'gang-arg', used for the 'gang' clause.
- bool ParseOpenACCGangArg(SourceLocation GangLoc);
+ bool ParseOpenACCGangArgList(SourceLocation GangLoc,
+ llvm::SmallVectorImpl<OpenACCGangKind> &GKs,
+ llvm::SmallVectorImpl<Expr *> &IntExprs);
+
+ using OpenACCGangArgRes = std::pair<OpenACCGangKind, ExprResult>;
+ /// Parses a 'gang-arg', used for the 'gang' clause. Returns a pair of the
+ /// ExprResult (which contains the validity of the expression), plus the gang
+ /// kind for the current argument.
+ OpenACCGangArgRes ParseOpenACCGangArg(SourceLocation GangLoc);
/// Parses a 'condition' expr, ensuring it results in a
ExprResult ParseOpenACCConditionExpr();
diff --git a/clang/include/clang/Sema/SemaOpenACC.h
b/clang/include/clang/Sema/SemaOpenACC.h
index 97386d2378b758..59a9648d5f9380 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -38,9 +38,20 @@ class SemaOpenACC : public SemaBase {
/// haven't had their 'parent' compute construct set yet. Entires will only
be
/// made to this list in the case where we know the loop isn't an orphan.
llvm::SmallVector<OpenACCLoopConstruct *> ParentlessLoopConstructs;
- /// Whether we are inside of a compute construct, and should add loops to the
- /// above collection.
- bool InsideComputeConstruct = false;
+
+ struct ComputeConstructInfo {
+ /// Which type of compute construct we are inside of, which we can use to
+ /// determine whether we should add loops to the above collection. We can
+ /// also use it to diagnose loop construct clauses.
+ OpenACCDirectiveKind Kind = OpenACCDirectiveKind::Invalid;
+ // If we have an active compute construct, stores the list of clauses we've
+ // prepared for it, so that we can diagnose limitations on child
constructs.
+ ArrayRef<OpenACCClause *> Clauses;
+ } ActiveComputeConstructInfo;
+
+ bool isInComputeConstruct() const {
+ return ActiveComputeConstructInfo.Kind != OpenACCDirectiveKind::Invalid;
+ }
/// 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
@@ -99,6 +110,15 @@ class SemaOpenACC : public SemaBase {
} TileInfo;
public:
+ ComputeConstructInfo &getActiveComputeConstructInfo() {
+ return ActiveComputeConstructInfo;
+ }
+
+ /// If there is a current 'active' loop construct with a 'gang' clause on a
+ /// 'kernel' construct, this will have the source location for it. This
+ /// permits us to implement the restriction of no further 'gang' clauses.
+ SourceLocation LoopGangClauseOnKernelLoc;
+
// Redeclaration of the version in OpenACCClause.h.
using DeviceTypeArgument = std::pair<IdentifierInfo *, SourceLocation>;
@@ -149,9 +169,14 @@ class SemaOpenACC : public SemaBase {
Expr *LoopCount;
};
+ struct GangDetails {
+ SmallVector<OpenACCGangKind> GangKinds;
+ SmallVector<Expr *> IntExprs;
+ };
+
std::variant<std::monostate, DefaultDetails, ConditionDetails,
IntExprDetails, VarListDetails, WaitDetails,
DeviceTypeDetails,
- ReductionDetails, CollapseDetails>
+ ReductionDetails, CollapseDetails, GangDetails>
Details = std::monostate{};
public:
@@ -245,9 +270,18 @@ class SemaOpenACC : public SemaBase {
ClauseKind == OpenACCClauseKind::NumWorkers ||
ClauseKind == OpenACCClauseKind::Async ||
ClauseKind == OpenACCClauseKind::Tile ||
+ ClauseKind == OpenACCClauseKind::Gang ||
ClauseKind == OpenACCClauseKind::VectorLength) &&
"Parsed clause kind does not have a int exprs");
+ if (ClauseKind == OpenACCClauseKind::Gang) {
+ // There might not be any gang int exprs, as this is an optional
+ // argument.
+ if (std::holds_alternative<std::monostate>(Details))
+ return {};
+ return std::get<GangDetails>(Details).IntExprs;
+ }
+
return std::get<IntExprDetails>(Details).IntExprs;
}
@@ -259,6 +293,16 @@ class SemaOpenACC : public SemaBase {
return std::get<ReductionDetails>(Details).Op;
}
+ ArrayRef<OpenACCGangKind> getGangKinds() const {
+ assert(ClauseKind == OpenACCClauseKind::Gang &&
+ "Parsed clause kind does not have gang kind");
+ // The args on gang are optional, so this might not actually hold
+ // anything.
+ if (std::holds_alternative<std::monostate>(Details))
+ return {};
+ return std::get<GangDetails>(Details).GangKinds;
+ }
+
ArrayRef<Expr *> getVarList() {
assert((ClauseKind == OpenACCClauseKind::Private ||
ClauseKind == OpenACCClauseKind::NoCreate ||
@@ -371,6 +415,25 @@ class SemaOpenACC : public SemaBase {
Details = IntExprDetails{std::move(IntExprs)};
}
+ void setGangDetails(ArrayRef<OpenACCGangKind> GKs,
+ ArrayRef<Expr *> IntExprs) {
+ assert(ClauseKind == OpenACCClauseKind::Gang &&
+ "Parsed Clause kind does not have gang details");
+ assert(GKs.size() == IntExprs.size() && "Mismatched kind/size?");
+
+ Details = GangDetails{{GKs.begin(), GKs.end()},
+ {IntExprs.begin(), IntExprs.end()}};
+ }
+
+ void setGangDetails(llvm::SmallVector<OpenACCGangKind> &&GKs,
+ llvm::SmallVector<Expr *> &&IntExprs) {
+ assert(ClauseKind == OpenACCClauseKind::Gang &&
+ "Parsed Clause kind does not have gang details");
+ assert(GKs.size() == IntExprs.size() && "Mismatched kind/size?");
+
+ Details = GangDetails{std::move(GKs), std::move(IntExprs)};
+ }
+
void setVarListDetails(ArrayRef<Expr *> VarList, bool IsReadOnly,
bool IsZero) {
assert((ClauseKind == OpenACCClauseKind::Private ||
@@ -545,10 +608,12 @@ class SemaOpenACC : public SemaBase {
SourceLocation RBLoc);
/// Checks the loop depth value for a collapse clause.
ExprResult CheckCollapseLoopCount(Expr *LoopCount);
- /// Checks a single size expr for a tile clause. 'gang' could possibly call
- /// this, but has slightly stricter rules as to valid values.
+ /// Checks a single size expr for a tile clause.
ExprResult CheckTileSizeExpr(Expr *SizeExpr);
+ // Check a single expression on a gang clause.
+ ExprResult CheckGangExpr(OpenACCGangKind GK, Expr *E);
+
ExprResult BuildOpenACCAsteriskSizeExpr(SourceLocation AsteriskLoc);
ExprResult ActOnOpenACCAsteriskSizeExpr(SourceLocation AsteriskLoc);
@@ -595,8 +660,9 @@ class SemaOpenACC : public SemaBase {
/// Loop needing its parent construct.
class AssociatedStmtRAII {
SemaOpenACC &SemaRef;
- bool WasInsideComputeConstruct;
+ ComputeConstructInfo OldActiveComputeConstructInfo;
OpenACCDirectiveKind DirKind;
+ SourceLocation OldLoopGangClauseOnKernelLoc;
llvm::SmallVector<OpenACCLoopConstruct *> ParentlessLoopConstructs;
LoopInConstructRAII LoopRAII;
diff --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp
index 0b34ed6189593e..6fb8fe0b8cfeef 100644
--- a/clang/lib/AST/OpenACCClause.cpp
+++ b/clang/lib/AST/OpenACCClause.cpp
@@ -26,7 +26,7 @@ bool OpenACCClauseWithExprs::classof(const OpenACCClause *C) {
return OpenACCWaitClause::classof(C) || OpenACCNumGangsClause::classof(C) ||
OpenACCTileClause::classof(C) ||
OpenACCClauseWithSingleIntExpr::classof(C) ||
- OpenACCClauseWithVarList::classof(C);
+ OpenACCGangClause::classof(C) || OpenACCClauseWithVarList::classof(C);
}
bool OpenACCClauseWithVarList::classof(const OpenACCClause *C) {
return OpenACCPrivateClause::classof(C) ||
@@ -125,6 +125,21 @@
OpenACCNumWorkersClause::OpenACCNumWorkersClause(SourceLocation BeginLoc,
"Condition expression type not scalar/dependent");
}
+OpenACCGangClause::OpenACCGangClause(SourceLocation BeginLoc,
+ SourceLocation LParenLoc,
+ ArrayRef<OpenACCGangKind> GangKinds,
+ ArrayRef<Expr *> IntExprs,
+ SourceLocation EndLoc)
+ : OpenACCClauseWithExprs(OpenACCClauseKind::Gang, BeginLoc, LParenLoc,
+ EndLoc) {
+ assert(GangKinds.size() == IntExprs.size() && "Mismatch exprs/kind?");
+ std::uninitialized_copy(IntExprs.begin(), IntExprs.end(),
+ getTrailingObjects<Expr *>());
+ setExprs(MutableArrayRef(getTrailingObjects<Expr *>(), IntExprs.size()));
+ std::uninitialized_copy(GangKinds.begin(), GangKinds.end(),
+ getTrailingObjects<OpenACCGangKind>());
+}
+
OpenACCNumWorkersClause *
OpenACCNumWorkersClause::Create(const ASTContext &C, SourceLocation BeginLoc,
SourceLocation LParenLoc, Expr *IntExpr,
@@ -376,11 +391,16 @@ OpenACCSeqClause *OpenACCSeqClause::Create(const
ASTContext &C,
return new (Mem) OpenACCSeqClause(BeginLoc, EndLoc);
}
-OpenACCGangClause *OpenACCGangClause::Create(const ASTContext &C,
- SourceLocation BeginLoc,
- SourceLocation EndLoc) {
- void *Mem = C.Allocate(sizeof(OpenACCGangClause));
- return new (Mem) OpenACCGangClause(BeginLoc, EndLoc);
+OpenACCGangClause *
+OpenACCGangClause::Create(const ASTContext &C, SourceLocation BeginLoc,
+ SourceLocation LParenLoc,
+ ArrayRef<OpenACCGangKind> GangKinds,
+ ArrayRef<Expr *> IntExprs, SourceLocation EndLoc) {
+ void *Mem =
+ C.Allocate(OpenACCGangClause::totalSizeToAlloc<Expr *, OpenACCGangKind>(
+ IntExprs.size(), GangKinds.size()));
+ return new (Mem)
+ OpenACCGangClause(BeginLoc, LParenLoc, GangKinds, IntExprs, EndLoc);
}
OpenACCWorkerClause *OpenACCWorkerClause::Create(const ASTContext &C,
@@ -600,3 +620,21 @@ void OpenACCClausePrinter::VisitCollapseClause(const
OpenACCCollapseClause &C) {
printExpr(C.getLoopCount());
OS << ")";
}
+
+void OpenACCClausePrinter::VisitGangClause(const OpenACCGangClause &C) {
+ OS << "gang";
+
+ if (C.getNumExprs() > 0) {
+ OS << "(";
+ bool first = true;
+ for (unsigned I = 0; I < C.getNumExprs(); ++I) {
+ if (!first)
+ OS << ", ";
+ first = false;
+
+ OS << C.getExpr(I).first << ": ";
+ printExpr(C.getExpr(I).second);
+ }
+ OS << ")";
+ }
+}
diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index 4d177fd6c5968c..6161b1403ed35d 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -2646,6 +2646,12 @@ void OpenACCClauseProfiler::VisitIndependentClause(
void OpenACCClauseProfiler::VisitSeqClause(const OpenACCSeqClause &Clause) {}
+void OpenACCClauseProfiler::VisitGangClause(const OpenACCGangClause &Clause) {
+ for (unsigned I = 0; I < Clause.getNumExprs(); ++I) {
+ Profiler.VisitStmt(Clause.getExpr(I).second);
+ }
+}
+
void OpenACCClauseProfiler::VisitReductionClause(
const OpenACCReductionClause &Clause) {
for (auto *E : Clause.getVarList())
diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp
index 15b23d60c3ffab..ac8c196777f9b8 100644
--- a/clang/lib/AST/TextNodeDumper.cpp
+++ b/clang/lib/AST/TextNodeDumper.cpp
@@ -425,6 +425,17 @@ void TextNodeDumper::Visit(const OpenACCClause *C) {
// but print 'clause' here so it is clear what is happening from the
dump.
OS << " clause";
break;
+ case OpenACCClauseKind::Gang: {
+ OS << " clause";
+ // print the list of all GangKinds, so that there is some sort of
+ // relationship to the expressions listed afterwards.
+ auto *GC = cast<OpenACCGangClause>(C);
+
+ for (unsigned I = 0; I < GC->getNumExprs(); ++I) {
+ OS << " " << GC->getExpr(I).first;
+ }
+ break;
+ }
case OpenACCClauseKind::Collapse:
OS << " clause";
if (cast<OpenACCCollapseClause>(C)->hasForce())
diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp
index b27e50b147f4a8..635039b724e6a0 100644
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@@ -797,23 +797,26 @@ bool Parser::ParseOpenACCSizeExprList(
/// [num:]int-expr
/// dim:int-expr
/// static:size-expr
-bool Parser::ParseOpenACCGangArg(SourceLocation GangLoc) {
+Parser::OpenACCGangArgRes Parser::ParseOpenACCGangArg(SourceLocation GangLoc) {
if (isOpenACCSpecialToken(OpenACCSpecialTokenKind::Static, getCurToken()) &&
NextToken().is(tok::colon)) {
// 'static' just takes a size-expr, which is an int-expr or an asterisk.
ConsumeToken();
ConsumeToken();
- return ParseOpenACCSizeExpr(OpenACCClauseKind::Gang).isInvalid();
+ ExprResult Res = ParseOpenACCSizeExpr(OpenACCClauseKind::Gang);
+ return {OpenACCGangKind::Static, Res};
}
if (isOpenACCSpecialToken(OpenACCSpecialTokenKind::Dim, getCurToken()) &&
NextToken().is(tok::colon)) {
ConsumeToken();
ConsumeToken();
- return ParseOpenACCIntExpr(OpenACCDirectiveKind::Invalid,
- OpenACCClauseKind::Gang, GangLoc)
- .first.isInvalid();
+ // Parse this as a const-expression, and we'll check its integer-ness/value
+ // in CheckGangExpr.
+ ExprResult Res =
+ getActions().CorrectDelayedTyposInExpr(ParseConstantExpression());
+ return {OpenACCGangKind::Dim, Res};
}
if (isOpenACCSpecialToken(OpenACCSpecialTokenKind::Num, getCurToken()) &&
@@ -822,27 +825,40 @@ bool Parser::ParseOpenACCGangArg(SourceLocation GangLoc) {
ConsumeToken();
// Fallthrough to the 'int-expr' handling for when 'num' is omitted.
}
+
// This is just the 'num' case where 'num' is optional.
- return ParseOpenACCIntExpr(OpenACCDirectiveKind::Invalid,
- OpenACCClauseKind::Gang, GangLoc)
- .first.isInvalid();
+ ExprResult Res = ParseOpenACCIntExpr(OpenACCDirectiveKind::Invalid,
+ OpenACCClauseKind::Gang, GangLoc)
+ .first;
+ return {OpenACCGangKind::Num, Res};
}
-bool Parser::ParseOpenACCGangArgList(SourceLocation GangLoc) {
- if (ParseOpenACCGangArg(GangLoc)) {
+bool Parser::ParseOpenACCGangArgList(
+ SourceLocation GangLoc, llvm::SmallVectorImpl<OpenACCGangKind> &GKs,
+ llvm::SmallVectorImpl<Expr *> &IntExprs) {
+
+ Parser::OpenACCGangArgRes Res = ParseOpenACCGangArg(GangLoc);
+ if (!Res.second.isUsable()) {
SkipUntil(tok::r_paren, tok::annot_pragma_openacc_end,
Parser::StopBeforeMatch);
- return false;
+ return true;
}
+ GKs.push_back(Res.first);
+ IntExprs.push_back(Res.second.get());
+
while (!getCurToken().isOneOf(tok::r_paren, tok::annot_pragma_openacc_end)) {
ExpectAndConsume(tok::comma);
- if (ParseOpenACCGangArg(GangLoc)) {
+ Res = ParseOpenACCGangArg(GangLoc);
+ if (!Res.second.isUsable()) {
SkipUntil(tok::r_paren, tok::annot_pragma_openacc_end,
Parser::StopBeforeMatch);
- return false;
+ return true;
}
+
+ GKs.push_back(Res.first);
+ IntExprs.push_back(Res.second.get());
}
return false;
}
@@ -1129,12 +1145,16 @@ Parser::OpenACCClauseParseResult
Parser::ParseOpenACCClauseParams(
}
break;
}
- case OpenACCClauseKind::Gang:
- if (ParseOpenACCGangArgList(ClauseLoc)) {
+ case OpenACCClauseKind::Gang: {
+ llvm::SmallVector<OpenACCGangKind> GKs;
+ llvm::SmallVector<Expr *> IntExprs;
+ if (ParseOpenACCGangArgList(ClauseLoc, GKs, IntExprs)) {
Parens.skipToEnd();
return OpenACCCanContinue();
}
+ ParsedClause.setGangDetails(std::move(GKs), std::move(IntExprs));
break;
+ }
case OpenACCClauseKind::Wait: {
OpenACCWaitParseInfo Info =
ParseOpenACCWaitArgument(ClauseLoc,
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 66f8029a2754b9..30d73d621db69b 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -366,6 +366,19 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind
DirectiveKind,
}
}
+ case OpenACCClauseKind::Gang: {
+ switch (DirectiveKind) {
+ case OpenACCDirectiveKind::Loop:
+ case OpenACCDirectiveKind::ParallelLoop:
+ case OpenACCDirectiveKind::SerialLoop:
+ case OpenACCDirectiveKind::KernelsLoop:
+ case OpenACCDirectiveKind::Routine:
+ return true;
+ default:
+ return false;
+ }
+ }
+
default:
// Do nothing so we can go to the 'unimplemented' diagnostic instead.
return true;
@@ -459,6 +472,23 @@ class SemaOpenACCClauseVisitor {
return nullptr;
}
+ // OpenACC 3.3 2.9:
+ // A 'gang', 'worker', or 'vector' clause may not appear if a 'seq' clause
+ // appears.
+ bool DiagIfSeqClause(SemaOpenACC::OpenACCParsedClause &Clause) {
+ const auto *Itr =
+ llvm::find_if(ExistingClauses, llvm::IsaPred<OpenACCSeqClause>);
+
+ if (Itr != ExistingClauses.end()) {
+ SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_cannot_combine)
+ << Clause.getClauseKind() << (*Itr)->getClauseKind();
+ SemaRef.Diag((*Itr)->getBeginLoc(), diag::note_acc_previous_clause_here);
+
+ return true;
+ }
+ return false;
+ }
+
public:
SemaOpenACCClauseVisitor(SemaOpenACC &S,
ArrayRef<const OpenACCClause *> ExistingClauses)
@@ -470,26 +500,14 @@ class SemaOpenACCClauseVisitor {
OpenACCClause *Visit(SemaOpenACC::OpenACCParsedClause &Clause) {
switch (Clause.getClauseKind()) {
- case OpenACCClauseKind::Gang:
- case OpenACCClauseKind::Worker:
- case OpenACCClauseKind::Vector: {
- // TODO OpenACC: These are only implemented enough for the 'seq'
diagnostic,
- // otherwise treats itself as unimplemented. When we implement these, we
- // can remove them from here.
-
- // OpenACC 3.3 2.9:
- // A 'gang', 'worker', or 'vector' clause may not appear if a 'seq' clause
- // appears.
- const auto *Itr =
- llvm::find_if(ExistingClauses, llvm::IsaPred<OpenACCSeqClause>);
-
- if (Itr != ExistingClauses.end()) {
- SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_cannot_combine)
- << Clause.getClauseKind() << (*Itr)->getClauseKind();
- SemaRef.Diag((*Itr)->getBeginLoc(), diag::note_acc_previous_clause_here);
+ case OpenACCClauseKind::Worker:
+ case OpenACCClauseKind::Vector: {
+ // TODO OpenACC: These are only implemented enough for the 'seq'
+ // diagnostic, otherwise treats itself as unimplemented. When we
+ // implement these, we can remove them from here.
+ DiagIfSeqClause(Clause);
+ return isNotImplemented();
}
- return isNotImplemented();
- }
#define VISIT_CLAUSE(CLAUSE_NAME)
\
case OpenACCClauseKind::CLAUSE_NAME:
\
@@ -1006,6 +1024,84 @@ OpenACCClause
*SemaOpenACCClauseVisitor::VisitIndependentClause(
Clause.getEndLoc());
}
+OpenACCClause *SemaOpenACCClauseVisitor::VisitGangClause(
+ SemaOpenACC::OpenACCParsedClause &Clause) {
+ if (DiagIfSeqClause(Clause))
+ return nullptr;
+
+ // Restrictions only properly implemented on 'loop' constructs, and it is
+ // the only construct that can do anything with this, so skip/treat as
+ // unimplemented for the combined constructs.
+ if (Clause.getDirectiveKind() != OpenACCDirectiveKind::Loop)
+ return isNotImplemented();
+
+ llvm::SmallVector<OpenACCGangKind> GangKinds;
+ llvm::SmallVector<Expr *> IntExprs;
+
+ // Store the existing locations, so we can do duplicate checking. Index is
+ // the int-value of the OpenACCGangKind enum.
+ SourceLocation ExistingElemLoc[3];
+
+ for (unsigned I = 0; I < Clause.getIntExprs().size(); ++I) {
+ OpenACCGangKind GK = Clause.getGangKinds()[I];
+ ExprResult ER = SemaRef.CheckGangExpr(GK, Clause.getIntExprs()[I]);
+
+ if (!ER.isUsable())
+ continue;
+
+ // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels
+ // construct, the gang clause behaves as follows. ... An argument with no
+ // keyword or with num keyword is only allowed when num_gangs does not
+ // appear on the kernels construct.
+ if (SemaRef.getActiveComputeConstructInfo().Kind ==
+ OpenACCDirectiveKind::Kernels &&
+ GK == OpenACCGangKind::Num) {
+ const auto *Itr =
+ llvm::find_if(SemaRef.getActiveComputeConstructInfo().Clauses,
+ llvm::IsaPred<OpenACCNumGangsClause>);
+
+ if (Itr != SemaRef.getActiveComputeConstructInfo().Clauses.end()) {
+ SemaRef.Diag(ER.get()->getBeginLoc(),
+ diag::err_acc_gang_num_gangs_conflict);
+ SemaRef.Diag((*Itr)->getBeginLoc(),
+ diag::note_acc_previous_clause_here);
+ continue;
+ }
+ }
+
+ // OpenACC 3.3 2.9: 'gang-arg-list' may have at most one num, one dim, and
+ // one static argument.
+ if (ExistingElemLoc[static_cast<unsigned>(GK)].isValid()) {
+ SemaRef.Diag(ER.get()->getBeginLoc(), diag::err_acc_gang_multiple_elt)
+ << static_cast<unsigned>(GK);
+ SemaRef.Diag(ExistingElemLoc[static_cast<unsigned>(GK)],
+ diag::note_acc_previous_expr_here);
+ continue;
+ }
+
+ ExistingElemLoc[static_cast<unsigned>(GK)] = ER.get()->getBeginLoc();
+ GangKinds.push_back(GK);
+ IntExprs.push_back(ER.get());
+ }
+
+ // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels
+ // construct, the gang clause behaves as follows. ... The region of a loop
+ // with a gang clause may not contain another loop with a gang clause unless
+ // within a nested compute region.
+ if (SemaRef.LoopGangClauseOnKernelLoc.isValid()) {
+ // This handles the 'inner loop' diagnostic, but we cannot set that we're
on
+ // one of these until we get to the end of the construct.
+ SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_gang_inside_gang);
+ SemaRef.Diag(SemaRef.LoopGangClauseOnKernelLoc,
+ diag::note_acc_previous_clause_here);
+ return nullptr;
+ }
+
+ return OpenACCGangClause::Create(Ctx, Clause.getBeginLoc(),
+ Clause.getLParenLoc(), GangKinds, IntExprs,
+ Clause.getEndLoc());
+}
+
OpenACCClause *SemaOpenACCClauseVisitor::VisitSeqClause(
SemaOpenACC::OpenACCParsedClause &Clause) {
// Restrictions only properly implemented on 'loop' constructs, and it is
@@ -1118,17 +1214,44 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
SemaOpenACC &S, OpenACCDirectiveKind DK,
ArrayRef<const OpenACCClause *> UnInstClauses,
ArrayRef<OpenACCClause *> Clauses)
- : SemaRef(S), WasInsideComputeConstruct(S.InsideComputeConstruct),
- DirKind(DK), LoopRAII(SemaRef, /*PreserveDepth=*/false) {
+ : SemaRef(S), OldActiveComputeConstructInfo(S.ActiveComputeConstructInfo),
+ DirKind(DK), OldLoopGangClauseOnKernelLoc(S.LoopGangClauseOnKernelLoc),
+ LoopRAII(SemaRef, /*PreserveDepth=*/false) {
// Compute constructs end up taking their 'loop'.
if (DirKind == OpenACCDirectiveKind::Parallel ||
DirKind == OpenACCDirectiveKind::Serial ||
DirKind == OpenACCDirectiveKind::Kernels) {
- SemaRef.InsideComputeConstruct = true;
+ SemaRef.ActiveComputeConstructInfo.Kind = DirKind;
+ SemaRef.ActiveComputeConstructInfo.Clauses = Clauses;
SemaRef.ParentlessLoopConstructs.swap(ParentlessLoopConstructs);
+
+ // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels
+ // construct, the gang clause behaves as follows. ... The region of a loop
+ // with a gang clause may not contain another loop with a gang clause
unless
+ // within a nested compute region.
+ //
+ // Implement the 'unless within a nested compute region' part.
+ SemaRef.LoopGangClauseOnKernelLoc = {};
} else if (DirKind == OpenACCDirectiveKind::Loop) {
SetCollapseInfoBeforeAssociatedStmt(UnInstClauses, Clauses);
SetTileInfoBeforeAssociatedStmt(UnInstClauses, Clauses);
+
+ // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels
+ // construct, the gang clause behaves as follows. ... The region of a loop
+ // with a gang clause may not contain another loop with a gang clause
unless
+ // within a nested compute region.
+ //
+ // We don't bother doing this when this is a template instantiation, as
+ // there is no reason to do these checks: the existance of a
+ // gang/kernels/etc cannot be dependent.
+ if (SemaRef.getActiveComputeConstructInfo().Kind ==
+ OpenACCDirectiveKind::Kernels &&
+ UnInstClauses.empty()) {
+ // This handles the 'outer loop' part of this.
+ auto *Itr = llvm::find_if(Clauses, llvm::IsaPred<OpenACCGangClause>);
+ if (Itr != Clauses.end())
+ SemaRef.LoopGangClauseOnKernelLoc = (*Itr)->getBeginLoc();
+ }
}
}
@@ -1199,7 +1322,9 @@ void
SemaOpenACC::AssociatedStmtRAII::SetTileInfoBeforeAssociatedStmt(
}
SemaOpenACC::AssociatedStmtRAII::~AssociatedStmtRAII() {
- SemaRef.InsideComputeConstruct = WasInsideComputeConstruct;
+ SemaRef.ActiveComputeConstructInfo = OldActiveComputeConstructInfo;
+ SemaRef.LoopGangClauseOnKernelLoc = OldLoopGangClauseOnKernelLoc;
+
if (DirKind == OpenACCDirectiveKind::Parallel ||
DirKind == OpenACCDirectiveKind::Serial ||
DirKind == OpenACCDirectiveKind::Kernels) {
@@ -1761,6 +1886,109 @@ ExprResult SemaOpenACC::CheckCollapseLoopCount(Expr
*LoopCount) {
ConstantExpr::Create(getASTContext(), LoopCount, APValue{*ICE})};
}
+namespace {
+ExprResult CheckGangStaticExpr(SemaOpenACC &S, Expr *E) {
+ if (isa<OpenACCAsteriskSizeExpr>(E))
+ return E;
+ return S.ActOnIntExpr(OpenACCDirectiveKind::Invalid, OpenACCClauseKind::Gang,
+ E->getBeginLoc(), E);
+}
+} // namespace
+
+ExprResult SemaOpenACC::CheckGangExpr(OpenACCGangKind GK, Expr *E) {
+ // Gang Expr legality depends on the associated compute construct.
+ switch (ActiveComputeConstructInfo.Kind) {
+ case OpenACCDirectiveKind::Invalid:
+ case OpenACCDirectiveKind::Parallel: {
+ switch (GK) {
+ // OpenACC 3.3 2.9.2: When the parent compute construct is a parallel
+ // construct, or an orphaned loop construct, the gang clause behaves as
+ // follows. ... The dim argument must be a constant positive integer
value
+ // 1, 2, or 3.
+ case OpenACCGangKind::Dim: {
+ if (!E)
+ return ExprError();
+ ExprResult Res =
+ ActOnIntExpr(OpenACCDirectiveKind::Invalid, OpenACCClauseKind::Gang,
+ E->getBeginLoc(), E);
+
+ if (!Res.isUsable())
+ return Res;
+
+ if (Res.get()->isInstantiationDependent())
+ return Res;
+
+ std::optional<llvm::APSInt> ICE =
+ Res.get()->getIntegerConstantExpr(getASTContext());
+
+ if (!ICE || *ICE <= 0 || ICE > 3) {
+ Diag(Res.get()->getBeginLoc(), diag::err_acc_gang_dim_value)
+ << ICE.has_value() << ICE.value_or(llvm::APSInt{}).getExtValue();
+ return ExprError();
+ }
+
+ return ExprResult{
+ ConstantExpr::Create(getASTContext(), Res.get(), APValue{*ICE})};
+ }
+ // OpenACC 3.3 2.9.2: When the parent compute construct is a parallel
+ // construct, or an orphaned loop construct, the gang clause behaves as
+ // follows. ... The num argument is not allowed.
+ case OpenACCGangKind::Num:
+ Diag(E->getBeginLoc(), diag::err_acc_gang_arg_invalid)
+ << GK
+ << (/*orphan/parallel=*/ActiveComputeConstructInfo.Kind ==
+ OpenACCDirectiveKind::Parallel
+ ? 1
+ : 0);
+ return ExprError();
+ case OpenACCGangKind::Static:
+ return CheckGangStaticExpr(*this, E);
+ }
+ } break;
+ case OpenACCDirectiveKind::Kernels: {
+ switch (GK) {
+ // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels
+ // construct, the gang clause behaves as follows. ... The dim argument is
+ // not allowed.
+ case OpenACCGangKind::Dim:
+ Diag(E->getBeginLoc(), diag::err_acc_gang_arg_invalid)
+ << GK << /*kernels=*/2;
+ return ExprError();
+ // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels
+ // construct, the gang clause behaves as follows. ... An argument with no
+ // keyword or with num keyword is only allowed when num_gangs does not
+ // appear on the kernels construct. ... The region of a loop with the gang
+ // clause may not contain another loop with a gang clause unless within a
+ // nested compute region.
+ case OpenACCGangKind::Num:
+ // This isn't allowed if there is a 'num_gangs' on the kernel construct,
+ // and makes loop-with-gang-clause ill-formed inside of this 'loop', but
+ // nothing can be enforced here.
+ return ExprResult{E};
+ case OpenACCGangKind::Static:
+ return CheckGangStaticExpr(*this, E);
+ }
+ } break;
+ case OpenACCDirectiveKind::Serial: {
+ switch (GK) {
+ // 'dim' and 'num' don't really make sense on serial, and GCC rejects them
+ // too, so we disallow them too.
+ case OpenACCGangKind::Dim:
+ case OpenACCGangKind::Num:
+ Diag(E->getBeginLoc(), diag::err_acc_gang_arg_invalid)
+ << GK << /*Kernels=*/3;
+ return ExprError();
+ case OpenACCGangKind::Static:
+ return CheckGangStaticExpr(*this, E);
+ }
+ }
+ default:
+ llvm_unreachable("Non compute construct in active compute construct?");
+ }
+
+ llvm_unreachable("Compute construct directive not handled?");
+}
+
ExprResult SemaOpenACC::CheckTileSizeExpr(Expr *SizeExpr) {
if (!SizeExpr)
return ExprError();
@@ -2031,7 +2259,7 @@ StmtResult
SemaOpenACC::ActOnEndStmtDirective(OpenACCDirectiveKind K,
// If we are in the scope of a compute construct, add this to the list of
// loop constructs that need assigning to the next closing compute
// construct.
- if (InsideComputeConstruct)
+ if (isInComputeConstruct())
ParentlessLoopConstructs.push_back(LoopConstruct);
return LoopConstruct;
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index 5753c9eccf6c92..cde40773336866 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -11940,6 +11940,29 @@ void OpenACCClauseTransform<Derived>::VisitTileClause(
ParsedClause.getLParenLoc(), ParsedClause.getIntExprs(),
ParsedClause.getEndLoc());
}
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitGangClause(
+ const OpenACCGangClause &C) {
+ llvm::SmallVector<OpenACCGangKind> TransformedGangKinds;
+ llvm::SmallVector<Expr *> TransformedIntExprs;
+
+ for (unsigned I = 0; I < C.getNumExprs(); ++I) {
+ ExprResult ER = Self.TransformExpr(const_cast<Expr
*>(C.getExpr(I).second));
+ if (!ER.isUsable())
+ continue;
+
+ ER = Self.getSema().OpenACC().CheckGangExpr(C.getExpr(I).first, ER.get());
+ if (!ER.isUsable())
+ continue;
+ TransformedGangKinds.push_back(C.getExpr(I).first);
+ TransformedIntExprs.push_back(ER.get());
+ }
+
+ NewClause = OpenACCGangClause::Create(
+ Self.getSema().getASTContext(), ParsedClause.getBeginLoc(),
+ ParsedClause.getLParenLoc(), TransformedGangKinds, TransformedIntExprs,
+ ParsedClause.getEndLoc());
+}
} // namespace
template <typename Derived>
OpenACCClause *TreeTransform<Derived>::TransformOpenACCClause(
diff --git a/clang/lib/Serialization/ASTReader.cpp
b/clang/lib/Serialization/ASTReader.cpp
index e638129897692f..0339419da43cab 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -12326,6 +12326,18 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
return OpenACCTileClause::Create(getContext(), BeginLoc, LParenLoc,
SizeExprs, EndLoc);
}
+ case OpenACCClauseKind::Gang: {
+ SourceLocation LParenLoc = readSourceLocation();
+ unsigned NumExprs = readInt();
+ llvm::SmallVector<OpenACCGangKind> GangKinds;
+ llvm::SmallVector<Expr *> Exprs;
+ for (unsigned I = 0; I < NumExprs; ++I) {
+ GangKinds.push_back(readEnum<OpenACCGangKind>());
+ Exprs.push_back(readSubExpr());
+ }
+ return OpenACCGangClause::Create(getContext(), BeginLoc, LParenLoc,
+ GangKinds, Exprs, EndLoc);
+ }
case OpenACCClauseKind::Finalize:
case OpenACCClauseKind::IfPresent:
@@ -12342,7 +12354,6 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
case OpenACCClauseKind::Bind:
case OpenACCClauseKind::DeviceNum:
case OpenACCClauseKind::DefaultAsync:
- case OpenACCClauseKind::Gang:
case OpenACCClauseKind::Invalid:
llvm_unreachable("Clause serialization not yet implemented");
}
diff --git a/clang/lib/Serialization/ASTWriter.cpp
b/clang/lib/Serialization/ASTWriter.cpp
index 4976327fc654ee..583d9a4bccb800 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -8182,6 +8182,16 @@ void ASTRecordWriter::writeOpenACCClause(const
OpenACCClause *C) {
AddStmt(E);
return;
}
+ case OpenACCClauseKind::Gang: {
+ const auto *GC = cast<OpenACCGangClause>(C);
+ writeSourceLocation(GC->getLParenLoc());
+ writeUInt32(GC->getNumExprs());
+ for (unsigned I = 0; I < GC->getNumExprs(); ++I) {
+ writeEnum(GC->getExpr(I).first);
+ AddStmt(const_cast<Expr *>(GC->getExpr(I).second));
+ }
+ return;
+ }
case OpenACCClauseKind::Finalize:
case OpenACCClauseKind::IfPresent:
@@ -8198,7 +8208,6 @@ void ASTRecordWriter::writeOpenACCClause(const
OpenACCClause *C) {
case OpenACCClauseKind::Bind:
case OpenACCClauseKind::DeviceNum:
case OpenACCClauseKind::DefaultAsync:
- case OpenACCClauseKind::Gang:
case OpenACCClauseKind::Invalid:
llvm_unreachable("Clause serialization not yet implemented");
}
diff --git a/clang/test/AST/ast-print-openacc-loop-construct.cpp
b/clang/test/AST/ast-print-openacc-loop-construct.cpp
index aee4591cab428f..baa4b173f88edc 100644
--- a/clang/test/AST/ast-print-openacc-loop-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-loop-construct.cpp
@@ -95,4 +95,86 @@ void foo() {
for(;;)
for(;;)
for(;;);
+
+// CHECK: #pragma acc loop gang(dim: 2)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc loop gang(dim:2)
+ for(;;);
+
+// CHECK: #pragma acc loop gang(static: i)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc loop gang(static:i)
+ for(;;);
+
+// CHECK: #pragma acc loop gang(static: i) gang(dim: 2)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc loop gang(static:i) gang(dim:2)
+ for(;;);
+
+// CHECK: #pragma acc parallel
+// CHECK-NEXT: #pragma acc loop gang(dim: 2)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc parallel
+#pragma acc loop gang(dim:2)
+ for(;;);
+
+// CHECK: #pragma acc parallel
+// CHECK-NEXT: #pragma acc loop gang(static: i)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc parallel
+#pragma acc loop gang(static:i)
+ for(;;);
+
+// CHECK: #pragma acc parallel
+// CHECK-NEXT: #pragma acc loop gang(static: i) gang(dim: 2)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc parallel
+#pragma acc loop gang(static:i) gang(dim:2)
+ for(;;);
+
+// CHECK: #pragma acc kernels
+// CHECK-NEXT: #pragma acc loop gang(num: i) gang(static: i)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc kernels
+#pragma acc loop gang(i) gang(static:i)
+ for(;;);
+
+// CHECK: #pragma acc kernels
+// CHECK-NEXT: #pragma acc loop gang(num: i) gang(static: i)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc kernels
+#pragma acc loop gang(num:i) gang(static:i)
+ for(;;);
+
+// CHECK: #pragma acc serial
+// CHECK-NEXT: #pragma acc loop gang(static: i)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc serial
+#pragma acc loop gang(static:i)
+ for(;;);
+
+// CHECK: #pragma acc serial
+// CHECK-NEXT: #pragma acc loop gang(static: *)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc serial
+#pragma acc loop gang(static:*)
+ for(;;);
+
+// CHECK: #pragma acc serial
+// CHECK-NEXT: #pragma acc loop
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc serial
+#pragma acc loop gang
+ for(;;);
}
diff --git a/clang/test/ParserOpenACC/parse-clauses.c
b/clang/test/ParserOpenACC/parse-clauses.c
index 6c382379a8a7ea..899fbd78b87298 100644
--- a/clang/test/ParserOpenACC/parse-clauses.c
+++ b/clang/test/ParserOpenACC/parse-clauses.c
@@ -1202,7 +1202,6 @@ void Tile() {
}
void Gang() {
- // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause
ignored}}
#pragma acc loop gang
for(;;){}
// expected-error@+3{{expected expression}}
@@ -1210,68 +1209,58 @@ void Gang() {
// expected-note@+1{{to match this '('}}
#pragma acc loop gang(
for(;;){}
- // expected-error@+2{{expected expression}}
- // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause
ignored}}
+ // expected-error@+1{{expected expression}}
#pragma acc loop gang()
for(;;){}
- // expected-error@+2{{expected expression}}
- // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause
ignored}}
+ // expected-error@+1{{expected expression}}
#pragma acc loop gang(5, *)
for(;;){}
- // expected-error@+2{{expected expression}}
- // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause
ignored}}
+ // expected-error@+1{{expected expression}}
#pragma acc loop gang(*)
for(;;){}
- // expected-error@+2{{expected expression}}
- // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause
ignored}}
+ // expected-error@+1{{expected expression}}
#pragma acc loop gang(5, num:*)
for(;;){}
- // expected-error@+2{{expected expression}}
- // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause
ignored}}
+ // expected-error@+1{{expected expression}}
#pragma acc loop gang(num:5, *)
for(;;){}
- // expected-error@+2{{expected expression}}
- // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause
ignored}}
+ // expected-error@+1{{expected expression}}
#pragma acc loop gang(num:5, num:*)
for(;;){}
- // expected-error@+2{{expected expression}}
- // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause
ignored}}
+ // expected-error@+1{{expected expression}}
#pragma acc loop gang(num:*)
for(;;){}
- // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause
ignored}}
-#pragma acc loop gang(dim:5)
+#pragma acc loop gang(dim:2)
for(;;){}
- // expected-error@+2{{expected expression}}
- // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause
ignored}}
+ // expected-error@+1{{expected expression}}
#pragma acc loop gang(dim:5, dim:*)
for(;;){}
- // expected-error@+2{{expected expression}}
- // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause
ignored}}
+ // expected-error@+1{{expected expression}}
#pragma acc loop gang(dim:*)
for(;;){}
- // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause
ignored}}
#pragma acc loop gang(static:*)
for(;;){}
- // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause
ignored}}
+ // expected-error@+2{{OpenACC 'gang' clause may have at most one 'static'
argument}}
+ // expected-note@+1{{previous expression is here}}
#pragma acc loop gang(static:*, static:5)
for(;;){}
- // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause
ignored}}
+#pragma acc kernels
#pragma acc loop gang(static:*, 5)
for(;;){}
- // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause
ignored}}
+#pragma acc kernels
#pragma acc loop gang(static:45, 5)
for(;;){}
@@ -1330,11 +1319,16 @@ void Gang() {
#pragma acc loop gang(dim:45
for(;;){}
- // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause
ignored}}
-#pragma acc loop gang(static:*, dim:returns_int(), 5)
+#pragma acc kernels
+#pragma acc loop gang(static:*, 5)
+ for(;;){}
+
+ // expected-error@+1{{argument to 'gang' clause dimension must be a constant
expression}}
+#pragma acc loop gang(static:*, dim:returns_int())
for(;;){}
- // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause
ignored}}
+ // expected-error@+2 2{{'num' argument on 'gang' clause is not permitted on
an orphaned 'loop' construct}}
+ // expected-error@+1{{argument to 'gang' clause dimension must be a constant
expression}}
#pragma acc loop gang(num: 32, static:*, dim:returns_int(), 5)
for(;;){}
diff --git a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
index d08497a7782edb..89000517c43fb5 100644
--- a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
@@ -214,8 +214,7 @@ void uses() {
// expected-error@+1{{OpenACC 'tile' clause is not valid on 'kernels'
directive}}
#pragma acc kernels device_type(*) tile(Var, 1)
while(1);
- // expected-error@+2{{OpenACC clause 'gang' may not follow a 'dtype' clause
in a compute construct}}
- // expected-note@+1{{previous clause is here}}
+ // expected-error@+1{{OpenACC 'gang' clause is not valid on 'kernels'
directive}}
#pragma acc kernels dtype(*) gang
while(1);
#pragma acc kernels device_type(*) wait
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 3da7f0e9836be8..6c2c79b02a4131 100644
--- a/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
+++ b/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
@@ -168,7 +168,6 @@ void uses() {
#pragma acc loop auto tile(1+2, 1)
for(;;)
for(;;);
- // expected-warning@+1{{OpenACC clause 'gang' not yet implemented}}
#pragma acc loop auto gang
for(;;);
// expected-error@+1{{OpenACC 'wait' clause is not valid on 'loop'
directive}}
@@ -306,7 +305,6 @@ void uses() {
#pragma acc loop tile(1+2, 1) auto
for(;;)
for(;;);
- // expected-warning@+1{{OpenACC clause 'gang' not yet implemented}}
#pragma acc loop gang auto
for(;;);
// expected-error@+1{{OpenACC 'wait' clause is not valid on 'loop'
directive}}
@@ -445,7 +443,6 @@ void uses() {
#pragma acc loop independent tile(1+2, 1)
for(;;)
for(;;);
- // expected-warning@+1{{OpenACC clause 'gang' not yet implemented}}
#pragma acc loop independent gang
for(;;);
// expected-error@+1{{OpenACC 'wait' clause is not valid on 'loop'
directive}}
@@ -583,7 +580,6 @@ void uses() {
#pragma acc loop tile(1+2, 1) independent
for(;;)
for(;;);
- // expected-warning@+1{{OpenACC clause 'gang' not yet implemented}}
#pragma acc loop gang independent
for(;;);
// expected-error@+1{{OpenACC 'wait' clause is not valid on 'loop'
directive}}
@@ -591,9 +587,8 @@ void uses() {
for(;;);
// 'seq' cannot be combined with 'gang', 'worker' or 'vector'
- // expected-error@+3{{OpenACC clause 'gang' may not appear on the same
construct as a 'seq' clause on a 'loop' construct}}
- // expected-note@+2{{previous clause is here}}
- // expected-warning@+1{{OpenACC clause 'gang' not yet implemented}}
+ // expected-error@+2{{OpenACC clause 'gang' may not appear on the same
construct as a 'seq' clause on a 'loop' construct}}
+ // expected-note@+1{{previous clause is here}}
#pragma acc loop seq gang
for(;;);
// expected-error@+3{{OpenACC clause 'worker' may not appear on the same
construct as a 'seq' clause on a 'loop' construct}}
@@ -735,10 +730,8 @@ void uses() {
#pragma acc loop seq wait
for(;;);
- // TODO OpenACC: when 'gang' is implemented and makes it to the AST, this
should diagnose because of a conflict with 'seq'.
- // TODOexpected-error@+3{{OpenACC clause 'gang' may not appear on the same
construct as a 'seq' clause on a 'loop' construct}}
- // TODOexpected-note@+2{{previous clause is here}}
- // expected-warning@+1{{OpenACC clause 'gang' not yet implemented}}
+ // expected-error@+2{{OpenACC clause 'seq' may not appear on the same
construct as a 'gang' clause on a 'loop' construct}}
+ // expected-note@+1{{previous clause is here}}
#pragma acc loop gang seq
for(;;);
// TODO OpenACC: when 'worker' is implemented and makes it to the AST, this
should diagnose because of a conflict with 'seq'.
diff --git a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
index 3d77c031f42630..cedef3ca858f5e 100644
--- a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
@@ -193,7 +193,6 @@ void uses() {
for(;;)
for(;;);
- // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause
ignored}}
#pragma acc loop dtype(*) gang
for(;;);
// expected-error@+1{{OpenACC 'wait' clause is not valid on 'loop'
directive}}
diff --git a/clang/test/SemaOpenACC/loop-construct-gang-ast.cpp
b/clang/test/SemaOpenACC/loop-construct-gang-ast.cpp
new file mode 100644
index 00000000000000..e797d842e240dc
--- /dev/null
+++ b/clang/test/SemaOpenACC/loop-construct-gang-ast.cpp
@@ -0,0 +1,330 @@
+// RUN: %clang_cc1 %s -fopenacc -ast-dump | FileCheck %s
+
+// Test this with PCH.
+// RUN: %clang_cc1 %s -fopenacc -emit-pch -o %t %s
+// RUN: %clang_cc1 %s -fopenacc -include-pch %t -ast-dump-all | FileCheck %s
+#ifndef PCH_HELPER
+#define PCH_HELPER
+
+void NormalUses() {
+ // CHECK: FunctionDecl{{.*}}NormalUses
+ // CHECK-NEXT: CompoundStmt
+
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: gang clause dim
+ // CHECK-NEXT: ConstantExpr{{.*}} 'int'
+ // CHECK-NEXT: value: Int 1
+ // CHECK-NEXT: IntegerLiteral{{.*}}'int' 1
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: NullStmt
+#pragma acc loop gang(dim:1)
+ for(;;);
+
+ int Val;
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} used Val 'int'
+
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: gang clause static
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'Val' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: NullStmt
+#pragma acc loop gang(static:Val)
+ for(;;);
+
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}}
kernels
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: gang clause num
+ // CHECK-NEXT: IntegerLiteral{{.*}}'int' 1
+ // CHECK-NEXT: gang clause static
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'Val' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: NullStmt
+#pragma acc kernels
+#pragma acc loop gang(num:1) gang(static:Val)
+ for(;;);
+
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}}
parallel
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: gang clause dim static
+ // CHECK-NEXT: ConstantExpr{{.*}} 'int'
+ // CHECK-NEXT: value: Int 1
+ // CHECK-NEXT: IntegerLiteral{{.*}}'int' 1
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'Val' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: NullStmt
+#pragma acc parallel
+#pragma acc loop gang(dim:1, static:Val)
+ for(;;);
+
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}}
serial
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: gang clause static
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'Val' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: NullStmt
+#pragma acc serial
+#pragma acc loop gang(static:Val)
+ for(;;);
+
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}}
serial
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: gang clause static
+ // CHECK-NEXT: OpenACCAsteriskSizeExpr
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: NullStmt
+#pragma acc serial
+#pragma acc loop gang(static:*)
+ for(;;);
+
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}}
serial
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: gang clause
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: NullStmt
+#pragma acc serial
+#pragma acc loop gang
+ for(;;);
+}
+
+template<typename T, unsigned One>
+void TemplateUses(T Val) {
+ // CHECK: FunctionTemplateDecl{{.*}}TemplateUses
+ // CHECK-NEXT: TemplateTypeParmDecl {{.*}} referenced typename depth 0 index
0 T
+ // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} referenced 'unsigned int' depth
0 index 1 One
+ // CHECK-NEXT: FunctionDecl{{.*}} TemplateUses 'void (T)'
+ // CHECK-NEXT: ParmVarDecl{{.*}} referenced Val 'T'
+ // CHECK-NEXT: CompoundStmt
+
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: gang clause dim
+ // CHECK-NEXT: DeclRefExpr{{.*}}'unsigned int' NonTypeTemplateParm{{.*}}
'One' 'unsigned int'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: NullStmt
+#pragma acc loop gang(dim:One)
+ for(;;);
+
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: gang clause static
+ // CHECK-NEXT: DeclRefExpr{{.*}}'T' lvalue ParmVar{{.*}} 'Val' 'T'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: NullStmt
+#pragma acc loop gang(static:Val)
+ for(;;);
+
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: gang clause static
+ // CHECK-NEXT: OpenACCAsteriskSizeExpr
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: NullStmt
+#pragma acc loop gang(static:*)
+ for(;;);
+
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}}
parallel
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: gang clause dim
+ // CHECK-NEXT: DeclRefExpr{{.*}}'unsigned int' NonTypeTemplateParm{{.*}}
'One' 'unsigned int'
+ // CHECK-NEXT: gang clause static
+ // CHECK-NEXT: DeclRefExpr{{.*}}'T' lvalue ParmVar{{.*}} 'Val' 'T'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: NullStmt
+#pragma acc parallel
+#pragma acc loop gang(dim:One) gang(static:Val)
+ for(;;);
+
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}}
parallel
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: gang clause dim static
+ // CHECK-NEXT: DeclRefExpr{{.*}}'unsigned int' NonTypeTemplateParm{{.*}}
'One' 'unsigned int'
+ // CHECK-NEXT: DeclRefExpr{{.*}}'T' lvalue ParmVar{{.*}} 'Val' 'T'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: NullStmt
+#pragma acc parallel
+#pragma acc loop gang(dim:One, static:Val)
+ for(;;);
+
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}}
serial
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: gang clause static
+ // CHECK-NEXT: DeclRefExpr{{.*}}'T' lvalue ParmVar{{.*}} 'Val' 'T'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: NullStmt
+#pragma acc serial
+#pragma acc loop gang(static:Val)
+ for(;;);
+
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}}
serial
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: gang clause
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: NullStmt
+#pragma acc serial
+#pragma acc loop gang
+ for(;;);
+
+ // Instantiation:
+ // CHECK-NEXT: FunctionDecl{{.*}} used TemplateUses 'void (int)'
implicit_instantiation
+ // CHECK-NEXT: TemplateArgument type 'int'
+ // CHECK-NEXT: BuiltinType{{.*}} 'int'
+ // CHECK-NEXT: TemplateArgument integral '1U'
+ // CHECK-NEXT: ParmVarDecl{{.*}} used Val 'int'
+ // CHECK-NEXT: CompoundStmt
+ //
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: gang clause dim
+ // CHECK-NEXT: ConstantExpr{{.*}} 'unsigned int'
+ // CHECK-NEXT: value: Int 1
+ // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}}'unsigned int'
+ // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'unsigned int' depth 0 index 1
One
+ // CHECK-NEXT: IntegerLiteral{{.*}}'unsigned int' 1
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: NullStmt
+ //
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: gang clause static
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+ // CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue ParmVar{{.*}} 'Val' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: NullStmt
+ //
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: gang clause static
+ // CHECK-NEXT: OpenACCAsteriskSizeExpr
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: NullStmt
+ //
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}}
parallel
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: gang clause dim
+ // CHECK-NEXT: ConstantExpr{{.*}} 'unsigned int'
+ // CHECK-NEXT: value: Int 1
+ // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}}'unsigned int'
+ // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'unsigned int' depth 0 index 1
One
+ // CHECK-NEXT: IntegerLiteral{{.*}}'unsigned int' 1
+ // CHECK-NEXT: gang clause static
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+ // CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue ParmVar{{.*}} 'Val' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: NullStmt
+ //
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}}
parallel
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: gang clause dim static
+ // CHECK-NEXT: ConstantExpr{{.*}} 'unsigned int'
+ // CHECK-NEXT: value: Int 1
+ // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}}'unsigned int'
+ // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'unsigned int' depth 0 index 1
One
+ // CHECK-NEXT: IntegerLiteral{{.*}}'unsigned int' 1
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+ // CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue ParmVar{{.*}} 'Val' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: NullStmt
+ //
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}}
serial
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: gang clause static
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+ // CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue ParmVar{{.*}} 'Val' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: NullStmt
+ //
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}}
serial
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: gang clause
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: <<<NULL>>
+ // CHECK-NEXT: NullStmt
+}
+
+void inst() {
+ TemplateUses<int, 1>(5);
+}
+
+#endif // PCH_HELPER
diff --git a/clang/test/SemaOpenACC/loop-construct-gang-clause.cpp
b/clang/test/SemaOpenACC/loop-construct-gang-clause.cpp
new file mode 100644
index 00000000000000..ab6439ae576193
--- /dev/null
+++ b/clang/test/SemaOpenACC/loop-construct-gang-clause.cpp
@@ -0,0 +1,335 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+struct S{};
+struct Converts{
+ operator int();
+};
+
+template<typename T, unsigned Zero, unsigned Two, unsigned Four>
+void ParallelOrOrphanTempl() {
+ T i;
+ // expected-error@+1{{'num' argument on 'gang' clause is not permitted on an
orphaned 'loop' construct}}
+#pragma acc loop gang(i)
+ for(;;);
+ // expected-error@+1{{'num' argument on 'gang' clause is not permitted on an
orphaned 'loop' construct}}
+#pragma acc loop gang(num:i)
+ for(;;);
+
+ // expected-error@+2{{'num' argument on 'gang' clause is not permitted on a
'loop' construct associated with a 'parallel' compute construct}}
+#pragma acc parallel
+#pragma acc loop gang(i)
+ for(;;);
+
+ // expected-error@+2{{'num' argument on 'gang' clause is not permitted on a
'loop' construct associated with a 'parallel' compute construct}}
+#pragma acc parallel
+#pragma acc loop gang(num:i)
+ for(;;);
+
+ // expected-error@+1{{argument to 'gang' clause dimension must be a constant
expression}}
+#pragma acc loop gang(dim:i)
+ for(;;);
+
+ // expected-error@+2{{argument to 'gang' clause dimension must be a constant
expression}}
+#pragma acc parallel
+#pragma acc loop gang(dim:i)
+ for(;;);
+
+ // expected-error@+1{{argument to 'gang' clause dimension must be 1, 2, or
3: evaluated to 0}}
+#pragma acc loop gang(dim:Zero)
+ for(;;);
+
+ // expected-error@+2{{argument to 'gang' clause dimension must be 1, 2, or
3: evaluated to 0}}
+#pragma acc parallel
+#pragma acc loop gang(dim:Zero)
+ for(;;);
+
+ // expected-error@+1{{argument to 'gang' clause dimension must be 1, 2, or
3: evaluated to 4}}
+#pragma acc loop gang(dim:Four)
+ for(;;);
+
+ // expected-error@+2{{argument to 'gang' clause dimension must be 1, 2, or
3: evaluated to 4}}
+#pragma acc parallel
+#pragma acc loop gang(dim:Four)
+ for(;;);
+
+#pragma acc loop gang(static:i) gang(dim:Two)
+ for(;;);
+
+#pragma acc parallel
+#pragma acc loop gang(dim:Two) gang(static:*)
+ for(;;);
+
+#pragma acc parallel
+#pragma acc loop gang(dim:Two, static:i)
+ for(;;);
+
+ // expected-error@+4{{OpenACC 'gang' clause may have at most one 'static'
argument}}
+ // expected-note@+3{{previous expression is here}}
+ // expected-error@+2{{OpenACC 'gang' clause may have at most one 'dim'
argument}}
+ // expected-note@+1{{previous expression is here}}
+#pragma acc loop gang(static:i, static:i, dim:Two, dim:1)
+ for(;;);
+}
+
+void ParallelOrOrphan() {
+ ParallelOrOrphanTempl<int, 0, 2, 4>(); // expected-note{{in instantiation of
function template}}
+
+ int i;
+ // expected-error@+1{{'num' argument on 'gang' clause is not permitted on an
orphaned 'loop' construct}}
+#pragma acc loop gang(i)
+ for(;;);
+ // expected-error@+1{{'num' argument on 'gang' clause is not permitted on an
orphaned 'loop' construct}}
+#pragma acc loop gang(num:i)
+ for(;;);
+
+ // expected-error@+2{{'num' argument on 'gang' clause is not permitted on a
'loop' construct associated with a 'parallel' compute construct}}
+#pragma acc parallel
+#pragma acc loop gang(i)
+ for(;;);
+
+ // expected-error@+2{{'num' argument on 'gang' clause is not permitted on a
'loop' construct associated with a 'parallel' compute construct}}
+#pragma acc parallel
+#pragma acc loop gang(num:i)
+ for(;;);
+
+ // expected-error@+1{{argument to 'gang' clause dimension must be a constant
expression}}
+#pragma acc loop gang(dim:i)
+ for(;;);
+
+ // expected-error@+2{{argument to 'gang' clause dimension must be a constant
expression}}
+#pragma acc parallel
+#pragma acc loop gang(dim:i)
+ for(;;);
+
+ // expected-error@+1{{argument to 'gang' clause dimension must be 1, 2, or
3: evaluated to 0}}
+#pragma acc loop gang(dim:0)
+ for(;;);
+
+ // expected-error@+2{{argument to 'gang' clause dimension must be 1, 2, or
3: evaluated to 0}}
+#pragma acc parallel
+#pragma acc loop gang(dim:0)
+ for(;;);
+
+ // expected-error@+1{{argument to 'gang' clause dimension must be 1, 2, or
3: evaluated to 4}}
+#pragma acc loop gang(dim:4)
+ for(;;);
+
+ // expected-error@+2{{argument to 'gang' clause dimension must be 1, 2, or
3: evaluated to 4}}
+#pragma acc parallel
+#pragma acc loop gang(dim:4)
+ for(;;);
+
+#pragma acc loop gang(static:i) gang(dim:2)
+ for(;;);
+
+#pragma acc parallel
+#pragma acc loop gang(dim:2) gang(static:i)
+ for(;;);
+
+ S s;
+ // expected-error@+2{{OpenACC clause 'gang' requires expression of integer
type ('S' invalid)}}
+#pragma acc parallel
+#pragma acc loop gang(dim:2) gang(static:s)
+ for(;;);
+
+ Converts C;
+#pragma acc parallel
+#pragma acc loop gang(dim:2) gang(static:C)
+ for(;;);
+}
+
+template<typename SomeS, typename SomeC, typename Int>
+void StaticIsIntegralTempl() {
+ SomeS s;
+ // expected-error@+2{{OpenACC clause 'gang' requires expression of integer
type ('S' invalid)}}
+#pragma acc parallel
+#pragma acc loop gang(dim:2) gang(static:s)
+ for(;;);
+
+ SomeC C;
+#pragma acc parallel
+#pragma acc loop gang(dim:2) gang(static:C)
+ for(;;);
+ Int i;
+#pragma acc parallel
+#pragma acc loop gang(dim:2) gang(static:i)
+ for(;;);
+
+#pragma acc parallel
+#pragma acc loop gang(dim:2) gang(static:*)
+ for(;;);
+}
+
+void StaticIsIntegral() {
+ StaticIsIntegralTempl<S, Converts, int>();// expected-note{{in instantiation
of function template}}
+
+ S s;
+ // expected-error@+2{{OpenACC clause 'gang' requires expression of integer
type ('S' invalid)}}
+#pragma acc parallel
+#pragma acc loop gang(dim:2) gang(static:s)
+ for(;;);
+
+ Converts C;
+#pragma acc parallel
+#pragma acc loop gang(dim:2) gang(static:C)
+ for(;;);
+}
+
+template<unsigned I>
+void SerialTempl() {
+ // expected-error@+2{{'num' argument on 'gang' clause is not permitted on a
'loop' construct associated with a 'serial' compute construct}}
+#pragma acc serial
+#pragma acc loop gang(I)
+ for(;;);
+
+ // expected-error@+2{{'num' argument on 'gang' clause is not permitted on a
'loop' construct associated with a 'serial' compute construct}}
+#pragma acc serial
+#pragma acc loop gang(num:I)
+ for(;;);
+
+ // expected-error@+2{{'dim' argument on 'gang' clause is not permitted on a
'loop' construct associated with a 'serial' compute construct}}
+#pragma acc serial
+#pragma acc loop gang(dim:I)
+ for(;;);
+
+#pragma acc serial
+#pragma acc loop gang(static:I)
+ for(;;);
+}
+
+void Serial() {
+ SerialTempl<2>();
+
+ // expected-error@+2{{'num' argument on 'gang' clause is not permitted on a
'loop' construct associated with a 'serial' compute construct}}
+#pragma acc serial
+#pragma acc loop gang(1)
+ for(;;);
+
+ // expected-error@+2{{'num' argument on 'gang' clause is not permitted on a
'loop' construct associated with a 'serial' compute construct}}
+#pragma acc serial
+#pragma acc loop gang(num:1)
+ for(;;);
+
+ // expected-error@+2{{'dim' argument on 'gang' clause is not permitted on a
'loop' construct associated with a 'serial' compute construct}}
+#pragma acc serial
+#pragma acc loop gang(dim:1)
+ for(;;);
+
+#pragma acc serial
+#pragma acc loop gang(static:1)
+ for(;;);
+
+ int i;
+
+#pragma acc serial
+#pragma acc loop gang(static:i)
+ for(;;);
+}
+
+template<typename T>
+void KernelsTempl() {
+ T t;
+ // expected-error@+2{{'dim' argument on 'gang' clause is not permitted on a
'loop' construct associated with a 'kernels' compute construct}}
+#pragma acc kernels
+#pragma acc loop gang(dim:t)
+ for(;;);
+
+#pragma acc kernels
+#pragma acc loop gang(static:t)
+ for(;;);
+
+ // expected-error@+3{{'num' argument to 'gang' clause not allowed on a
'loop' construct associated with a 'kernels' construct that has a 'num_gangs'
clause}}
+ // expected-note@+1{{previous clause is here}}
+#pragma acc kernels num_gangs(t)
+#pragma acc loop gang(t)
+ for(;;);
+
+ // expected-error@+3{{'num' argument to 'gang' clause not allowed on a
'loop' construct associated with a 'kernels' construct that has a 'num_gangs'
clause}}
+ // expected-note@+1{{previous clause is here}}
+#pragma acc kernels num_gangs(t)
+#pragma acc loop gang(num:t)
+ for(;;);
+}
+
+void Kernels() {
+ KernelsTempl<unsigned>();
+
+ // expected-error@+2{{'dim' argument on 'gang' clause is not permitted on a
'loop' construct associated with a 'kernels' compute construct}}
+#pragma acc kernels
+#pragma acc loop gang(dim:1)
+ for(;;);
+ unsigned t;
+#pragma acc kernels
+#pragma acc loop gang(static:t)
+ for(;;);
+
+ // expected-error@+3{{'num' argument to 'gang' clause not allowed on a
'loop' construct associated with a 'kernels' construct that has a 'num_gangs'
clause}}
+ // expected-note@+1{{previous clause is here}}
+#pragma acc kernels num_gangs(1)
+#pragma acc loop gang(1)
+ for(;;);
+
+ // expected-error@+3{{'num' argument to 'gang' clause not allowed on a
'loop' construct associated with a 'kernels' construct that has a 'num_gangs'
clause}}
+ // expected-note@+1{{previous clause is here}}
+#pragma acc kernels num_gangs(1)
+#pragma acc loop gang(num:1)
+ for(;;);
+
+#pragma acc kernels
+#pragma acc loop gang(num:1)
+ for(;;) {
+ // expected-error@+2{{loop with a 'gang' clause may not exist in the
region of a 'gang' clause on a 'kernels' compute construct}}
+ // expected-note@-3{{previous clause is here}}
+#pragma acc loop gang(static:1)
+ for(;;);
+ }
+
+#pragma acc kernels
+#pragma acc loop gang(num:1)
+ for(;;) {
+ // allowed, intervening compute construct
+#pragma acc serial
+#pragma acc loop gang(static:1)
+ for(;;);
+ }
+
+#pragma acc kernels
+#pragma acc loop gang(num:1)
+ for(;;);
+
+ // OK, on a different 'loop', not in the assoc statement.
+#pragma acc loop gang(static:1)
+ for(;;);
+
+ // expected-error@+3{{OpenACC 'gang' clause may have at most one unnamed or
'num' argument}}
+ // expected-note@+2{{previous expression is here}}
+#pragma acc kernels
+#pragma acc loop gang(5, num:1)
+ for(;;);
+
+ // expected-error@+3{{OpenACC 'gang' clause may have at most one unnamed or
'num' argument}}
+ // expected-note@+2{{previous expression is here}}
+#pragma acc kernels
+#pragma acc loop gang(num:5, 1)
+ for(;;);
+
+ // expected-error@+3{{OpenACC 'gang' clause may have at most one unnamed or
'num' argument}}
+ // expected-note@+2{{previous expression is here}}
+#pragma acc kernels
+#pragma acc loop gang(num:5, num:1)
+ for(;;);
+}
+
+void MaxOneEntry() {
+ // expected-error@+3{{OpenACC 'gang' clause may have at most one 'static'
argument}}
+ // expected-note@+2{{previous expression is here}}
+#pragma acc kernels
+#pragma acc loop gang(static: 1, static:1)
+ for(;;);
+
+#pragma acc kernels
+#pragma acc loop gang gang(static:1)
+ for(;;);
+}
+
+
diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
index c282a9071391e7..2ffe47fbd74476 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -2907,6 +2907,11 @@ void OpenACCClauseEnqueue::VisitSeqClause(const
OpenACCSeqClause &C) {}
void OpenACCClauseEnqueue::VisitCollapseClause(const OpenACCCollapseClause &C)
{
Visitor.AddStmt(C.getLoopCount());
}
+void OpenACCClauseEnqueue::VisitGangClause(const OpenACCGangClause &C) {
+ for (unsigned I = 0; I < C.getNumExprs(); ++I) {
+ Visitor.AddStmt(C.getExpr(I).second);
+ }
+}
} // namespace
void EnqueueVisitor::EnqueueChildren(const OpenACCClause *C) {
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
