[clang] 5b25c31 - [OpenACC] Implement loop 'gang' clause. (#112006)
via cfe-commits
cfe-commits at lists.llvm.org
Fri Oct 11 09:05:23 PDT 2024
Author: Erich Keane
Date: 2024-10-11T09:05:19-07:00
New Revision: 5b25c31351ad1b10a3819411379b3258869c1e1b
URL: https://github.com/llvm/llvm-project/commit/5b25c31351ad1b10a3819411379b3258869c1e1b
DIFF: https://github.com/llvm/llvm-project/commit/5b25c31351ad1b10a3819411379b3258869c1e1b.diff
LOG: [OpenACC] Implement loop 'gang' clause. (#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.
Added:
clang/test/SemaOpenACC/loop-construct-gang-ast.cpp
clang/test/SemaOpenACC/loop-construct-gang-clause.cpp
Modified:
clang/include/clang/AST/OpenACCClause.h
clang/include/clang/Basic/DiagnosticSemaKinds.td
clang/include/clang/Basic/OpenACCClauses.def
clang/include/clang/Basic/OpenACCKinds.h
clang/include/clang/Parse/Parser.h
clang/include/clang/Sema/SemaOpenACC.h
clang/lib/AST/OpenACCClause.cpp
clang/lib/AST/StmtProfile.cpp
clang/lib/AST/TextNodeDumper.cpp
clang/lib/Parse/ParseOpenACC.cpp
clang/lib/Sema/SemaOpenACC.cpp
clang/lib/Sema/TreeTransform.h
clang/lib/Serialization/ASTReader.cpp
clang/lib/Serialization/ASTWriter.cpp
clang/test/AST/ast-print-openacc-loop-construct.cpp
clang/test/ParserOpenACC/parse-clauses.c
clang/test/SemaOpenACC/compute-construct-device_type-clause.c
clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
clang/test/SemaOpenACC/loop-construct-device_type-clause.c
clang/tools/libclang/CIndex.cpp
Removed:
################################################################################
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 at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
#pragma acc loop gang
for(;;){}
// expected-error at +3{{expected expression}}
@@ -1210,68 +1209,58 @@ void Gang() {
// expected-note at +1{{to match this '('}}
#pragma acc loop gang(
for(;;){}
- // expected-error at +2{{expected expression}}
- // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
+ // expected-error at +1{{expected expression}}
#pragma acc loop gang()
for(;;){}
- // expected-error at +2{{expected expression}}
- // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
+ // expected-error at +1{{expected expression}}
#pragma acc loop gang(5, *)
for(;;){}
- // expected-error at +2{{expected expression}}
- // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
+ // expected-error at +1{{expected expression}}
#pragma acc loop gang(*)
for(;;){}
- // expected-error at +2{{expected expression}}
- // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
+ // expected-error at +1{{expected expression}}
#pragma acc loop gang(5, num:*)
for(;;){}
- // expected-error at +2{{expected expression}}
- // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
+ // expected-error at +1{{expected expression}}
#pragma acc loop gang(num:5, *)
for(;;){}
- // expected-error at +2{{expected expression}}
- // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
+ // expected-error at +1{{expected expression}}
#pragma acc loop gang(num:5, num:*)
for(;;){}
- // expected-error at +2{{expected expression}}
- // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
+ // expected-error at +1{{expected expression}}
#pragma acc loop gang(num:*)
for(;;){}
- // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
-#pragma acc loop gang(dim:5)
+#pragma acc loop gang(dim:2)
for(;;){}
- // expected-error at +2{{expected expression}}
- // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
+ // expected-error at +1{{expected expression}}
#pragma acc loop gang(dim:5, dim:*)
for(;;){}
- // expected-error at +2{{expected expression}}
- // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
+ // expected-error at +1{{expected expression}}
#pragma acc loop gang(dim:*)
for(;;){}
- // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
#pragma acc loop gang(static:*)
for(;;){}
- // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
+ // expected-error at +2{{OpenACC 'gang' clause may have at most one 'static' argument}}
+ // expected-note at +1{{previous expression is here}}
#pragma acc loop gang(static:*, static:5)
for(;;){}
- // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
+#pragma acc kernels
#pragma acc loop gang(static:*, 5)
for(;;){}
- // expected-warning at +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 at +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 at +1{{argument to 'gang' clause dimension must be a constant expression}}
+#pragma acc loop gang(static:*, dim:returns_int())
for(;;){}
- // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
+ // expected-error at +2 2{{'num' argument on 'gang' clause is not permitted on an orphaned 'loop' construct}}
+ // expected-error at +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 at +1{{OpenACC 'tile' clause is not valid on 'kernels' directive}}
#pragma acc kernels device_type(*) tile(Var, 1)
while(1);
- // expected-error at +2{{OpenACC clause 'gang' may not follow a 'dtype' clause in a compute construct}}
- // expected-note at +1{{previous clause is here}}
+ // expected-error at +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 at +1{{OpenACC clause 'gang' not yet implemented}}
#pragma acc loop auto gang
for(;;);
// expected-error at +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 at +1{{OpenACC clause 'gang' not yet implemented}}
#pragma acc loop gang auto
for(;;);
// expected-error at +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 at +1{{OpenACC clause 'gang' not yet implemented}}
#pragma acc loop independent gang
for(;;);
// expected-error at +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 at +1{{OpenACC clause 'gang' not yet implemented}}
#pragma acc loop gang independent
for(;;);
// expected-error at +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 at +3{{OpenACC clause 'gang' may not appear on the same construct as a 'seq' clause on a 'loop' construct}}
- // expected-note at +2{{previous clause is here}}
- // expected-warning at +1{{OpenACC clause 'gang' not yet implemented}}
+ // expected-error at +2{{OpenACC clause 'gang' may not appear on the same construct as a 'seq' clause on a 'loop' construct}}
+ // expected-note at +1{{previous clause is here}}
#pragma acc loop seq gang
for(;;);
// expected-error at +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 at +3{{OpenACC clause 'gang' may not appear on the same construct as a 'seq' clause on a 'loop' construct}}
- // TODOexpected-note at +2{{previous clause is here}}
- // expected-warning at +1{{OpenACC clause 'gang' not yet implemented}}
+ // expected-error at +2{{OpenACC clause 'seq' may not appear on the same construct as a 'gang' clause on a 'loop' construct}}
+ // expected-note at +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 at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
#pragma acc loop dtype(*) gang
for(;;);
// expected-error at +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 at +1{{'num' argument on 'gang' clause is not permitted on an orphaned 'loop' construct}}
+#pragma acc loop gang(i)
+ for(;;);
+ // expected-error at +1{{'num' argument on 'gang' clause is not permitted on an orphaned 'loop' construct}}
+#pragma acc loop gang(num:i)
+ for(;;);
+
+ // expected-error at +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 at +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 at +1{{argument to 'gang' clause dimension must be a constant expression}}
+#pragma acc loop gang(dim:i)
+ for(;;);
+
+ // expected-error at +2{{argument to 'gang' clause dimension must be a constant expression}}
+#pragma acc parallel
+#pragma acc loop gang(dim:i)
+ for(;;);
+
+ // expected-error at +1{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 0}}
+#pragma acc loop gang(dim:Zero)
+ for(;;);
+
+ // expected-error at +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 at +1{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 4}}
+#pragma acc loop gang(dim:Four)
+ for(;;);
+
+ // expected-error at +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 at +4{{OpenACC 'gang' clause may have at most one 'static' argument}}
+ // expected-note at +3{{previous expression is here}}
+ // expected-error at +2{{OpenACC 'gang' clause may have at most one 'dim' argument}}
+ // expected-note at +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 at +1{{'num' argument on 'gang' clause is not permitted on an orphaned 'loop' construct}}
+#pragma acc loop gang(i)
+ for(;;);
+ // expected-error at +1{{'num' argument on 'gang' clause is not permitted on an orphaned 'loop' construct}}
+#pragma acc loop gang(num:i)
+ for(;;);
+
+ // expected-error at +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 at +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 at +1{{argument to 'gang' clause dimension must be a constant expression}}
+#pragma acc loop gang(dim:i)
+ for(;;);
+
+ // expected-error at +2{{argument to 'gang' clause dimension must be a constant expression}}
+#pragma acc parallel
+#pragma acc loop gang(dim:i)
+ for(;;);
+
+ // expected-error at +1{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 0}}
+#pragma acc loop gang(dim:0)
+ for(;;);
+
+ // expected-error at +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 at +1{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 4}}
+#pragma acc loop gang(dim:4)
+ for(;;);
+
+ // expected-error at +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 at +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 at +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 at +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 at +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 at +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 at +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 at +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 at +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 at +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 at +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 at +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 at +1{{previous clause is here}}
+#pragma acc kernels num_gangs(t)
+#pragma acc loop gang(t)
+ for(;;);
+
+ // expected-error at +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 at +1{{previous clause is here}}
+#pragma acc kernels num_gangs(t)
+#pragma acc loop gang(num:t)
+ for(;;);
+}
+
+void Kernels() {
+ KernelsTempl<unsigned>();
+
+ // expected-error at +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 at +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 at +1{{previous clause is here}}
+#pragma acc kernels num_gangs(1)
+#pragma acc loop gang(1)
+ for(;;);
+
+ // expected-error at +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 at +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 at +2{{loop with a 'gang' clause may not exist in the region of a 'gang' clause on a 'kernels' compute construct}}
+ // expected-note at -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
diff erent 'loop', not in the assoc statement.
+#pragma acc loop gang(static:1)
+ for(;;);
+
+ // expected-error at +3{{OpenACC 'gang' clause may have at most one unnamed or 'num' argument}}
+ // expected-note at +2{{previous expression is here}}
+#pragma acc kernels
+#pragma acc loop gang(5, num:1)
+ for(;;);
+
+ // expected-error at +3{{OpenACC 'gang' clause may have at most one unnamed or 'num' argument}}
+ // expected-note at +2{{previous expression is here}}
+#pragma acc kernels
+#pragma acc loop gang(num:5, 1)
+ for(;;);
+
+ // expected-error at +3{{OpenACC 'gang' clause may have at most one unnamed or 'num' argument}}
+ // expected-note at +2{{previous expression is here}}
+#pragma acc kernels
+#pragma acc loop gang(num:5, num:1)
+ for(;;);
+}
+
+void MaxOneEntry() {
+ // expected-error at +3{{OpenACC 'gang' clause may have at most one 'static' argument}}
+ // expected-note at +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) {
More information about the cfe-commits
mailing list