[clang] [OpenACC] Implement 'num_gangs' sema for compute constructs (PR #89460)
Erich Keane via cfe-commits
cfe-commits at lists.llvm.org
Fri Apr 19 14:45:17 PDT 2024
https://github.com/erichkeane created https://github.com/llvm/llvm-project/pull/89460
num_gangs takes an 'int-expr-list', for 'parallel', and an 'int-expr' for 'kernels'. This patch changes the parsing to always parse it as an 'int-expr-list', then correct the expression count during Sema. It also implements the rest of the semantic analysis changes for this clause.
>From aa7844c73b72172707ec81234c8e9d5370bbb772 Mon Sep 17 00:00:00 2001
From: erichkeane <ekeane at nvidia.com>
Date: Fri, 19 Apr 2024 10:20:40 -0700
Subject: [PATCH] [OpenACC] Implement 'num_gangs' sema for compute constructs
num_gangs takes an 'int-expr-list', for 'parallel', and an 'int-expr' for
'kernels'. This patch changes the parsing to always parse it as an
'int-expr-list', then correct the expression count during Sema. It also
implements the rest of the semantic analysis changes for this clause.
---
clang/include/clang/AST/OpenACCClause.h | 87 ++++++++--
.../clang/Basic/DiagnosticSemaKinds.td | 5 +
clang/include/clang/Basic/OpenACCClauses.def | 1 +
clang/include/clang/Parse/Parser.h | 16 +-
clang/include/clang/Sema/SemaOpenACC.h | 16 +-
clang/lib/AST/OpenACCClause.cpp | 16 ++
clang/lib/AST/StmtProfile.cpp | 6 +
clang/lib/AST/TextNodeDumper.cpp | 1 +
clang/lib/Parse/ParseOpenACC.cpp | 87 ++++++++--
clang/lib/Sema/SemaOpenACC.cpp | 35 ++++
clang/lib/Sema/TreeTransform.h | 26 +++
clang/lib/Serialization/ASTReader.cpp | 10 +-
clang/lib/Serialization/ASTWriter.cpp | 9 +-
clang/test/ParserOpenACC/parse-clauses.c | 4 -
.../compute-construct-intexpr-clause-ast.cpp | 79 +++++++++
.../compute-construct-num_gangs-clause.c | 54 ++++++
.../compute-construct-num_gangs-clause.cpp | 160 ++++++++++++++++++
clang/tools/libclang/CIndex.cpp | 4 +
18 files changed, 572 insertions(+), 44 deletions(-)
create mode 100644 clang/test/SemaOpenACC/compute-construct-num_gangs-clause.c
create mode 100644 clang/test/SemaOpenACC/compute-construct-num_gangs-clause.cpp
diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h
index 8b6d3221aa066b..277a351c49fcb8 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -156,33 +156,88 @@ class OpenACCSelfClause : public OpenACCClauseWithCondition {
Expr *ConditionExpr, SourceLocation EndLoc);
};
-/// Represents oen of a handful of classes that have a single integer
+/// Represents a clause that has one or more IntExprs. It does not own the
+/// IntExprs, but provides 'children' and other accessors.
+class OpenACCClauseWithIntExprs : public OpenACCClauseWithParams {
+ MutableArrayRef<Expr *> IntExprs;
+
+protected:
+ OpenACCClauseWithIntExprs(OpenACCClauseKind K, SourceLocation BeginLoc,
+ SourceLocation LParenLoc, SourceLocation EndLoc)
+ : OpenACCClauseWithParams(K, BeginLoc, LParenLoc, EndLoc) {}
+
+ /// Used only for initialization, the leaf class can initialize this to
+ /// trailing storage.
+ void setIntExprs(MutableArrayRef<Expr *> NewIntExprs) {
+ assert(IntExprs.empty() && "Cannot change IntExprs list");
+ IntExprs = NewIntExprs;
+ }
+
+ /// Gets the entire list of integer expressions, but leave it to the
+ /// individual clauses to expose this how they'd like.
+ llvm::ArrayRef<Expr *> getIntExprs() const { return IntExprs; }
+
+public:
+ child_range children() {
+ return child_range(reinterpret_cast<Stmt **>(IntExprs.begin()),
+ reinterpret_cast<Stmt **>(IntExprs.end()));
+ }
+
+ const_child_range children() const {
+ child_range Children =
+ const_cast<OpenACCClauseWithIntExprs *>(this)->children();
+ return const_child_range(Children.begin(), Children.end());
+ }
+};
+
+class OpenACCNumGangsClause final
+ : public OpenACCClauseWithIntExprs,
+ public llvm::TrailingObjects<OpenACCNumGangsClause, Expr *> {
+
+ OpenACCNumGangsClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
+ ArrayRef<Expr *> IntExprs, SourceLocation EndLoc)
+ : OpenACCClauseWithIntExprs(OpenACCClauseKind::NumGangs, BeginLoc,
+ LParenLoc, EndLoc) {
+ std::uninitialized_copy(IntExprs.begin(), IntExprs.end(),
+ getTrailingObjects<Expr *>());
+ setIntExprs(MutableArrayRef(getTrailingObjects<Expr *>(), IntExprs.size()));
+ }
+
+public:
+ static OpenACCNumGangsClause *
+ Create(const ASTContext &C, SourceLocation BeginLoc, SourceLocation LParenLoc,
+ ArrayRef<Expr *> IntExprs, SourceLocation EndLoc);
+
+ llvm::ArrayRef<Expr *> getIntExprs() {
+ return OpenACCClauseWithIntExprs::getIntExprs();
+ }
+
+ llvm::ArrayRef<Expr *> getIntExprs() const {
+ return OpenACCClauseWithIntExprs::getIntExprs();
+ }
+};
+
+/// Represents one of a handful of clauses that have a single integer
/// expression.
-class OpenACCClauseWithSingleIntExpr : public OpenACCClauseWithParams {
+class OpenACCClauseWithSingleIntExpr : public OpenACCClauseWithIntExprs {
Expr *IntExpr;
protected:
OpenACCClauseWithSingleIntExpr(OpenACCClauseKind K, SourceLocation BeginLoc,
SourceLocation LParenLoc, Expr *IntExpr,
SourceLocation EndLoc)
- : OpenACCClauseWithParams(K, BeginLoc, LParenLoc, EndLoc),
- IntExpr(IntExpr) {}
+ : OpenACCClauseWithIntExprs(K, BeginLoc, LParenLoc, EndLoc),
+ IntExpr(IntExpr) {
+ setIntExprs(MutableArrayRef<Expr *>{&this->IntExpr, 1});
+ }
public:
- bool hasIntExpr() const { return IntExpr; }
- const Expr *getIntExpr() const { return IntExpr; }
-
- Expr *getIntExpr() { return IntExpr; };
-
- child_range children() {
- return child_range(reinterpret_cast<Stmt **>(&IntExpr),
- reinterpret_cast<Stmt **>(&IntExpr + 1));
+ bool hasIntExpr() const { return !getIntExprs().empty(); }
+ const Expr *getIntExpr() const {
+ return hasIntExpr() ? getIntExprs()[0] : nullptr;
}
- const_child_range children() const {
- return const_child_range(reinterpret_cast<Stmt *const *>(&IntExpr),
- reinterpret_cast<Stmt *const *>(&IntExpr + 1));
- }
+ Expr *getIntExpr() { return hasIntExpr() ? getIntExprs()[0] : nullptr; };
};
class OpenACCNumWorkersClause : public OpenACCClauseWithSingleIntExpr {
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 1bbe76ff6bd2ac..a95424862e63f4 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12293,4 +12293,9 @@ def note_acc_int_expr_conversion
: Note<"conversion to %select{integral|enumeration}0 type %1">;
def err_acc_int_expr_multiple_conversions
: Error<"multiple conversions from expression type %0 to an integral type">;
+def err_acc_num_gangs_num_args
+ : Error<"%select{no|too many}0 integer expression arguments provided to "
+ "OpenACC 'num_gangs' "
+ "%select{|clause: '%1' directive expects maximum of %2, %3 were "
+ "provided}0">;
} // end of sema component.
diff --git a/clang/include/clang/Basic/OpenACCClauses.def b/clang/include/clang/Basic/OpenACCClauses.def
index 520e068f4ffd40..dd5792e7ca8c39 100644
--- a/clang/include/clang/Basic/OpenACCClauses.def
+++ b/clang/include/clang/Basic/OpenACCClauses.def
@@ -18,6 +18,7 @@
VISIT_CLAUSE(Default)
VISIT_CLAUSE(If)
VISIT_CLAUSE(Self)
+VISIT_CLAUSE(NumGangs)
VISIT_CLAUSE(NumWorkers)
VISIT_CLAUSE(VectorLength)
diff --git a/clang/include/clang/Parse/Parser.h b/clang/include/clang/Parse/Parser.h
index 72b2f958a5e622..7e95a52e6f34d5 100644
--- a/clang/include/clang/Parse/Parser.h
+++ b/clang/include/clang/Parse/Parser.h
@@ -3644,10 +3644,22 @@ class Parser : public CodeCompletionHandler {
/// Parses the clause of the 'bind' argument, which can be a string literal or
/// an ID expression.
ExprResult ParseOpenACCBindClauseArgument();
+
+ /// A type to represent the state of parsing after an attempt to parse an
+ /// OpenACC int-expr. This is useful to determine whether an int-expr list can
+ /// continue parsing after a failed int-expr.
+ using OpenACCIntExprParseResult =
+ std::pair<ExprResult, OpenACCParseCanContinue>;
/// Parses the clause kind of 'int-expr', which can be any integral
/// expression.
- ExprResult ParseOpenACCIntExpr(OpenACCDirectiveKind DK, OpenACCClauseKind CK,
- SourceLocation Loc);
+ OpenACCIntExprParseResult ParseOpenACCIntExpr(OpenACCDirectiveKind DK,
+ OpenACCClauseKind CK,
+ SourceLocation Loc);
+ /// Parses the argument list for 'num_gangs', which allows up to 3
+ /// 'int-expr's.
+ bool ParseOpenACCIntExprList(OpenACCDirectiveKind DK, OpenACCClauseKind CK,
+ SourceLocation Loc,
+ llvm::SmallVector<Expr *> &IntExprs);
/// Parses the 'device-type-list', which is a list of identifiers.
bool ParseOpenACCDeviceTypeList();
/// Parses the 'async-argument', which is an integral value with two
diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h
index 023722049732af..ea28617f79b81b 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -93,14 +93,16 @@ class SemaOpenACC : public SemaBase {
}
unsigned getNumIntExprs() const {
- assert((ClauseKind == OpenACCClauseKind::NumWorkers ||
+ assert((ClauseKind == OpenACCClauseKind::NumGangs ||
+ ClauseKind == OpenACCClauseKind::NumWorkers ||
ClauseKind == OpenACCClauseKind::VectorLength) &&
"Parsed clause kind does not have a int exprs");
return std::get<IntExprDetails>(Details).IntExprs.size();
}
ArrayRef<Expr *> getIntExprs() {
- assert((ClauseKind == OpenACCClauseKind::NumWorkers ||
+ assert((ClauseKind == OpenACCClauseKind::NumGangs ||
+ ClauseKind == OpenACCClauseKind::NumWorkers ||
ClauseKind == OpenACCClauseKind::VectorLength) &&
"Parsed clause kind does not have a int exprs");
return std::get<IntExprDetails>(Details).IntExprs;
@@ -134,11 +136,19 @@ class SemaOpenACC : public SemaBase {
}
void setIntExprDetails(ArrayRef<Expr *> IntExprs) {
- assert((ClauseKind == OpenACCClauseKind::NumWorkers ||
+ assert((ClauseKind == OpenACCClauseKind::NumGangs ||
+ ClauseKind == OpenACCClauseKind::NumWorkers ||
ClauseKind == OpenACCClauseKind::VectorLength) &&
"Parsed clause kind does not have a int exprs");
Details = IntExprDetails{{IntExprs.begin(), IntExprs.end()}};
}
+ void setIntExprDetails(llvm::SmallVector<Expr *> &&IntExprs) {
+ assert((ClauseKind == OpenACCClauseKind::NumGangs ||
+ ClauseKind == OpenACCClauseKind::NumWorkers ||
+ ClauseKind == OpenACCClauseKind::VectorLength) &&
+ "Parsed clause kind does not have a int exprs");
+ Details = IntExprDetails{IntExprs};
+ }
};
SemaOpenACC(Sema &S);
diff --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp
index 75334223e073c3..6cd5b28802187d 100644
--- a/clang/lib/AST/OpenACCClause.cpp
+++ b/clang/lib/AST/OpenACCClause.cpp
@@ -124,6 +124,16 @@ OpenACCVectorLengthClause::Create(const ASTContext &C, SourceLocation BeginLoc,
OpenACCVectorLengthClause(BeginLoc, LParenLoc, IntExpr, EndLoc);
}
+OpenACCNumGangsClause *OpenACCNumGangsClause::Create(const ASTContext &C,
+ SourceLocation BeginLoc,
+ SourceLocation LParenLoc,
+ ArrayRef<Expr *> IntExprs,
+ SourceLocation EndLoc) {
+ void *Mem = C.Allocate(
+ OpenACCNumGangsClause::totalSizeToAlloc<Expr *>(IntExprs.size()));
+ return new (Mem) OpenACCNumGangsClause(BeginLoc, LParenLoc, IntExprs, EndLoc);
+}
+
//===----------------------------------------------------------------------===//
// OpenACC clauses printing methods
//===----------------------------------------------------------------------===//
@@ -141,6 +151,12 @@ void OpenACCClausePrinter::VisitSelfClause(const OpenACCSelfClause &C) {
OS << "(" << CondExpr << ")";
}
+void OpenACCClausePrinter::VisitNumGangsClause(const OpenACCNumGangsClause &C) {
+ OS << "num_gangs(";
+ llvm::interleaveComma(C.getIntExprs(), OS);
+ OS << ")";
+}
+
void OpenACCClausePrinter::VisitNumWorkersClause(
const OpenACCNumWorkersClause &C) {
OS << "num_workers(" << C.getIntExpr() << ")";
diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index 8138ae3a244a8b..c81724f84dd9ce 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -2497,6 +2497,12 @@ void OpenACCClauseProfiler::VisitSelfClause(const OpenACCSelfClause &Clause) {
Profiler.VisitStmt(Clause.getConditionExpr());
}
+void OpenACCClauseProfiler::VisitNumGangsClause(
+ const OpenACCNumGangsClause &Clause) {
+ for (auto *E : Clause.getIntExprs())
+ Profiler.VisitStmt(E);
+}
+
void OpenACCClauseProfiler::VisitNumWorkersClause(
const OpenACCNumWorkersClause &Clause) {
assert(Clause.hasIntExpr() && "num_workers clause requires a valid int expr");
diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp
index e5a8b285715b30..8f0a9a9b0ed0bc 100644
--- a/clang/lib/AST/TextNodeDumper.cpp
+++ b/clang/lib/AST/TextNodeDumper.cpp
@@ -399,6 +399,7 @@ void TextNodeDumper::Visit(const OpenACCClause *C) {
break;
case OpenACCClauseKind::If:
case OpenACCClauseKind::Self:
+ case OpenACCClauseKind::NumGangs:
case OpenACCClauseKind::NumWorkers:
case OpenACCClauseKind::VectorLength:
// The condition expression will be printed as a part of the 'children',
diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp
index 757417f75c9636..9b7cd1668f7c17 100644
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@@ -632,16 +632,54 @@ Parser::ParseOpenACCClauseList(OpenACCDirectiveKind DirKind) {
return Clauses;
}
-ExprResult Parser::ParseOpenACCIntExpr(OpenACCDirectiveKind DK,
- OpenACCClauseKind CK,
- SourceLocation Loc) {
- ExprResult ER =
- getActions().CorrectDelayedTyposInExpr(ParseAssignmentExpression());
+Parser::OpenACCIntExprParseResult
+Parser::ParseOpenACCIntExpr(OpenACCDirectiveKind DK, OpenACCClauseKind CK,
+ SourceLocation Loc) {
+ ExprResult ER = ParseAssignmentExpression();
+ // If the actual parsing failed, we don't know the state of the parse, so
+ // don't try to continue.
if (!ER.isUsable())
- return ER;
+ return {ER, OpenACCParseCanContinue::Cannot};
+
+ // Parsing can continue after the initial assignment expression parsing, so
+ // even if there was a typo, we can continue.
+ ER = getActions().CorrectDelayedTyposInExpr(ER);
+ if (!ER.isUsable())
+ return {ER, OpenACCParseCanContinue::Can};
+
+ return {getActions().OpenACC().ActOnIntExpr(DK, CK, Loc, ER.get()),
+ OpenACCParseCanContinue::Can};
+}
+
+bool Parser::ParseOpenACCIntExprList(OpenACCDirectiveKind DK,
+ OpenACCClauseKind CK, SourceLocation Loc,
+ llvm::SmallVector<Expr *> &IntExprs) {
+ OpenACCIntExprParseResult CurResult = ParseOpenACCIntExpr(DK, CK, Loc);
+
+ if (!CurResult.first.isUsable() &&
+ CurResult.second == OpenACCParseCanContinue::Cannot) {
+ SkipUntil(tok::r_paren, tok::annot_pragma_openacc_end,
+ Parser::StopBeforeMatch);
+ return true;
+ }
+
+ IntExprs.push_back(CurResult.first.get());
+
+ while (!getCurToken().isOneOf(tok::r_paren, tok::annot_pragma_openacc_end)) {
+ ExpectAndConsume(tok::comma);
+
+ CurResult = ParseOpenACCIntExpr(DK, CK, Loc);
- return getActions().OpenACC().ActOnIntExpr(DK, CK, Loc, ER.get());
+ if (!CurResult.first.isUsable() &&
+ CurResult.second == OpenACCParseCanContinue::Cannot) {
+ SkipUntil(tok::r_paren, tok::annot_pragma_openacc_end,
+ Parser::StopBeforeMatch);
+ return true;
+ }
+ IntExprs.push_back(CurResult.first.get());
+ }
+ return false;
}
bool Parser::ParseOpenACCClauseVarList(OpenACCClauseKind Kind) {
@@ -761,7 +799,7 @@ bool Parser::ParseOpenACCGangArg(SourceLocation GangLoc) {
ConsumeToken();
return ParseOpenACCIntExpr(OpenACCDirectiveKind::Invalid,
OpenACCClauseKind::Gang, GangLoc)
- .isInvalid();
+ .first.isInvalid();
}
if (isOpenACCSpecialToken(OpenACCSpecialTokenKind::Num, getCurToken()) &&
@@ -773,7 +811,7 @@ bool Parser::ParseOpenACCGangArg(SourceLocation GangLoc) {
// This is just the 'num' case where 'num' is optional.
return ParseOpenACCIntExpr(OpenACCDirectiveKind::Invalid,
OpenACCClauseKind::Gang, GangLoc)
- .isInvalid();
+ .first.isInvalid();
}
bool Parser::ParseOpenACCGangArgList(SourceLocation GangLoc) {
@@ -946,13 +984,25 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
}
break;
}
- case OpenACCClauseKind::NumGangs:
+ case OpenACCClauseKind::NumGangs: {
+ llvm::SmallVector<Expr *> IntExprs;
+
+ if (ParseOpenACCIntExprList(OpenACCDirectiveKind::Invalid,
+ OpenACCClauseKind::NumGangs, ClauseLoc,
+ IntExprs)) {
+ Parens.skipToEnd();
+ return OpenACCCanContinue();
+ }
+ ParsedClause.setIntExprDetails(std::move(IntExprs));
+ break;
+ }
case OpenACCClauseKind::NumWorkers:
case OpenACCClauseKind::DeviceNum:
case OpenACCClauseKind::DefaultAsync:
case OpenACCClauseKind::VectorLength: {
ExprResult IntExpr = ParseOpenACCIntExpr(OpenACCDirectiveKind::Invalid,
- ClauseKind, ClauseLoc);
+ ClauseKind, ClauseLoc)
+ .first;
if (IntExpr.isInvalid()) {
Parens.skipToEnd();
return OpenACCCanContinue();
@@ -1017,7 +1067,8 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
: OpenACCSpecialTokenKind::Num,
ClauseKind);
ExprResult IntExpr = ParseOpenACCIntExpr(OpenACCDirectiveKind::Invalid,
- ClauseKind, ClauseLoc);
+ ClauseKind, ClauseLoc)
+ .first;
if (IntExpr.isInvalid()) {
Parens.skipToEnd();
return OpenACCCanContinue();
@@ -1081,11 +1132,13 @@ bool Parser::ParseOpenACCWaitArgument(SourceLocation Loc, bool IsDirective) {
// Consume colon.
ConsumeToken();
- ExprResult IntExpr = ParseOpenACCIntExpr(
- IsDirective ? OpenACCDirectiveKind::Wait
- : OpenACCDirectiveKind::Invalid,
- IsDirective ? OpenACCClauseKind::Invalid : OpenACCClauseKind::Wait,
- Loc);
+ ExprResult IntExpr =
+ ParseOpenACCIntExpr(IsDirective ? OpenACCDirectiveKind::Wait
+ : OpenACCDirectiveKind::Invalid,
+ IsDirective ? OpenACCClauseKind::Invalid
+ : OpenACCClauseKind::Wait,
+ Loc)
+ .first;
if (IntExpr.isInvalid())
return true;
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 190739fa02e932..ba69e71e30a181 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -91,6 +91,7 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
default:
return false;
}
+ case OpenACCClauseKind::NumGangs:
case OpenACCClauseKind::NumWorkers:
case OpenACCClauseKind::VectorLength:
switch (DirectiveKind) {
@@ -230,6 +231,40 @@ SemaOpenACC::ActOnClause(ArrayRef<const OpenACCClause *> ExistingClauses,
getASTContext(), Clause.getBeginLoc(), Clause.getLParenLoc(),
Clause.getConditionExpr(), Clause.getEndLoc());
}
+ case OpenACCClauseKind::NumGangs: {
+ // Restrictions only properly implemented on 'compute' constructs, and
+ // 'compute' constructs are the only construct that can do anything with
+ // this yet, so skip/treat as unimplemented in this case.
+ if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()))
+ break;
+
+ // There is no prose in the standard that says duplicates aren't allowed,
+ // but this diagnostic is present in other compilers, as well as makes
+ // sense.
+ if (checkAlreadyHasClauseOfKind(*this, ExistingClauses, Clause))
+ return nullptr;
+
+ if (Clause.getIntExprs().empty())
+ Diag(Clause.getBeginLoc(), diag::err_acc_num_gangs_num_args)
+ << /*NoArgs=*/0;
+
+ unsigned MaxArgs =
+ (Clause.getDirectiveKind() == OpenACCDirectiveKind::Parallel ||
+ Clause.getDirectiveKind() == OpenACCDirectiveKind::ParallelLoop)
+ ? 3
+ : 1;
+ if (Clause.getIntExprs().size() > MaxArgs)
+ Diag(Clause.getBeginLoc(), diag::err_acc_num_gangs_num_args)
+ << /*NoArgs=*/1 << Clause.getDirectiveKind() << MaxArgs
+ << Clause.getIntExprs().size();
+
+ // Create the AST node for the clause even if the number of expressions is
+ // incorrect.
+ return OpenACCNumGangsClause::Create(
+ getASTContext(), Clause.getBeginLoc(), Clause.getLParenLoc(),
+ Clause.getIntExprs(), Clause.getEndLoc());
+ break;
+ }
case OpenACCClauseKind::NumWorkers: {
// Restrictions only properly implemented on 'compute' constructs, and
// 'compute' constructs are the only construct that can do anything with
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index ade33ec65038fd..10ce226f1298dd 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -11159,6 +11159,32 @@ void OpenACCClauseTransform<Derived>::VisitSelfClause(
ParsedClause.getEndLoc());
}
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitNumGangsClause(
+ const OpenACCNumGangsClause &C) {
+ llvm::SmallVector<Expr *> InstantiatedIntExprs;
+
+ for (Expr *CurIntExpr : C.getIntExprs()) {
+ ExprResult Res = Self.TransformExpr(CurIntExpr);
+
+ if (!Res.isUsable())
+ return;
+
+ Res = Self.getSema().OpenACC().ActOnIntExpr(OpenACCDirectiveKind::Invalid,
+ C.getClauseKind(),
+ C.getBeginLoc(), Res.get());
+ if (!Res.isUsable())
+ return;
+
+ InstantiatedIntExprs.push_back(Res.get());
+ }
+
+ ParsedClause.setIntExprDetails(InstantiatedIntExprs);
+ NewClause = OpenACCNumGangsClause::Create(
+ Self.getSema().getASTContext(), ParsedClause.getBeginLoc(),
+ ParsedClause.getLParenLoc(), ParsedClause.getIntExprs(),
+ ParsedClause.getEndLoc());
+}
template <typename Derived>
void OpenACCClauseTransform<Derived>::VisitNumWorkersClause(
const OpenACCNumWorkersClause &C) {
diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
index 44e23919ea18e0..d64925676df7b1 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -11786,6 +11786,15 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
return OpenACCSelfClause::Create(getContext(), BeginLoc, LParenLoc,
CondExpr, EndLoc);
}
+ case OpenACCClauseKind::NumGangs: {
+ SourceLocation LParenLoc = readSourceLocation();
+ unsigned NumClauses = readInt();
+ llvm::SmallVector<Expr *> IntExprs;
+ for (unsigned I = 0; I < NumClauses; ++I)
+ IntExprs.push_back(readSubExpr());
+ return OpenACCNumGangsClause::Create(getContext(), BeginLoc, LParenLoc,
+ IntExprs, EndLoc);
+ }
case OpenACCClauseKind::NumWorkers: {
SourceLocation LParenLoc = readSourceLocation();
Expr *IntExpr = readSubExpr();
@@ -11826,7 +11835,6 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
case OpenACCClauseKind::Reduction:
case OpenACCClauseKind::Collapse:
case OpenACCClauseKind::Bind:
- case OpenACCClauseKind::NumGangs:
case OpenACCClauseKind::DeviceNum:
case OpenACCClauseKind::DefaultAsync:
case OpenACCClauseKind::DeviceType:
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index 8a4b36207c4734..018b854652a46e 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -7657,6 +7657,14 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
AddStmt(const_cast<Expr *>(SC->getConditionExpr()));
return;
}
+ case OpenACCClauseKind::NumGangs: {
+ const auto *NGC = cast<OpenACCNumGangsClause>(C);
+ writeSourceLocation(NGC->getLParenLoc());
+ writeUInt32(NGC->getIntExprs().size());
+ for (Expr *E : NGC->getIntExprs())
+ AddStmt(E);
+ return;
+ }
case OpenACCClauseKind::NumWorkers: {
const auto *NWC = cast<OpenACCNumWorkersClause>(C);
writeSourceLocation(NWC->getLParenLoc());
@@ -7697,7 +7705,6 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
case OpenACCClauseKind::Reduction:
case OpenACCClauseKind::Collapse:
case OpenACCClauseKind::Bind:
- case OpenACCClauseKind::NumGangs:
case OpenACCClauseKind::DeviceNum:
case OpenACCClauseKind::DefaultAsync:
case OpenACCClauseKind::DeviceType:
diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c
index ddf40f71701eda..799f22b8c120e5 100644
--- a/clang/test/ParserOpenACC/parse-clauses.c
+++ b/clang/test/ParserOpenACC/parse-clauses.c
@@ -911,16 +911,12 @@ void IntExprParsing() {
#pragma acc parallel num_gangs(invalid)
{}
- // expected-error at +2{{expected ')'}}
- // expected-note at +1{{to match this '('}}
#pragma acc parallel num_gangs(5, 4)
{}
- // expected-warning at +1{{OpenACC clause 'num_gangs' not yet implemented, clause ignored}}
#pragma acc parallel num_gangs(5)
{}
- // expected-warning at +1{{OpenACC clause 'num_gangs' not yet implemented, clause ignored}}
#pragma acc parallel num_gangs(returns_int())
{}
diff --git a/clang/test/SemaOpenACC/compute-construct-intexpr-clause-ast.cpp b/clang/test/SemaOpenACC/compute-construct-intexpr-clause-ast.cpp
index 889025c26818d8..5a4c9f05ee089e 100644
--- a/clang/test/SemaOpenACC/compute-construct-intexpr-clause-ast.cpp
+++ b/clang/test/SemaOpenACC/compute-construct-intexpr-clause-ast.cpp
@@ -88,6 +88,34 @@ void NormalUses() {
// CHECK-NEXT: WhileStmt
// CHECK-NEXT: CXXBoolLiteralExpr
// CHECK-NEXT: CompoundStmt
+
+#pragma acc parallel num_gangs(some_int(), some_long(), some_short())
+ while(true){}
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
+ // CHECK-NEXT: num_gangs clause
+ // CHECK-NEXT: CallExpr{{.*}}'int'
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'int (*)()' <FunctionToPointerDecay>
+ // CHECK-NEXT: DeclRefExpr{{.*}}'int ()' lvalue Function{{.*}} 'some_int' 'int ()'
+ // CHECK-NEXT: CallExpr{{.*}}'long'
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'long (*)()' <FunctionToPointerDecay>
+ // CHECK-NEXT: DeclRefExpr{{.*}}'long ()' lvalue Function{{.*}} 'some_long' 'long ()'
+ // CHECK-NEXT: CallExpr{{.*}}'short'
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'short (*)()' <FunctionToPointerDecay>
+ // CHECK-NEXT: DeclRefExpr{{.*}}'short ()' lvalue Function{{.*}} 'some_short' 'short ()'
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: CompoundStmt
+
+#pragma acc kernels num_gangs(some_int())
+ while(true){}
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels
+ // CHECK-NEXT: num_gangs clause
+ // CHECK-NEXT: CallExpr{{.*}}'int'
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'int (*)()' <FunctionToPointerDecay>
+ // CHECK-NEXT: DeclRefExpr{{.*}}'int ()' lvalue Function{{.*}} 'some_int' 'int ()'
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: CompoundStmt
}
template<typename T, typename U>
@@ -187,6 +215,31 @@ void TemplUses(T t, U u) {
// CHECK-NEXT: CXXBoolLiteralExpr
// CHECK-NEXT: CompoundStmt
+#pragma acc kernels num_gangs(u)
+ while(true){}
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels
+ // CHECK-NEXT: num_gangs clause
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'U' lvalue ParmVar{{.*}} 'u' 'U'
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: CompoundStmt
+
+#pragma acc parallel num_gangs(u, U::value)
+ while(true){}
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
+ // CHECK-NEXT: num_gangs clause
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'U' lvalue ParmVar{{.*}} 'u' 'U'
+ // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue
+ // CHECK-NEXT: NestedNameSpecifier TypeSpec 'U'
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: CompoundStmt
+
+
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}}EndMarker
+ int EndMarker;
+
// Check the instantiated versions of the above.
// CHECK-NEXT: FunctionDecl{{.*}} used TemplUses 'void (CorrectConvert, HasInt)' implicit_instantiation
// CHECK-NEXT: TemplateArgument type 'CorrectConvert'
@@ -288,6 +341,32 @@ void TemplUses(T t, U u) {
// CHECK-NEXT: WhileStmt
// CHECK-NEXT: CXXBoolLiteralExpr
// CHECK-NEXT: CompoundStmt
+
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels
+ // CHECK-NEXT: num_gangs clause
+ // CHECK-NEXT: ImplicitCastExpr{{.*}} 'char' <UserDefinedConversion>
+ // CHECK-NEXT: CXXMemberCallExpr{{.*}}'char'
+ // CHECK-NEXT: MemberExpr{{.*}} '<bound member function type>' .operator char
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'HasInt' lvalue ParmVar
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: CompoundStmt
+
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
+ // CHECK-NEXT: num_gangs clause
+ // CHECK-NEXT: ImplicitCastExpr{{.*}} 'char' <UserDefinedConversion>
+ // CHECK-NEXT: CXXMemberCallExpr{{.*}}'char'
+ // CHECK-NEXT: MemberExpr{{.*}} '<bound member function type>' .operator char
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'HasInt' lvalue ParmVar
+ // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'const int' lvalue Var{{.*}} 'value' 'const int'
+ // CHECK-NEXT: NestedNameSpecifier TypeSpec 'HasInt'
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: CompoundStmt
+
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}}EndMarker
}
struct HasInt {
diff --git a/clang/test/SemaOpenACC/compute-construct-num_gangs-clause.c b/clang/test/SemaOpenACC/compute-construct-num_gangs-clause.c
new file mode 100644
index 00000000000000..cdc6847b47f948
--- /dev/null
+++ b/clang/test/SemaOpenACC/compute-construct-num_gangs-clause.c
@@ -0,0 +1,54 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+short getS();
+void Test() {
+#pragma acc kernels num_gangs(1)
+ while(1);
+
+ // expected-error at +1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}}
+#pragma acc serial num_gangs(1)
+ while(1);
+
+#pragma acc parallel num_gangs(1)
+ while(1);
+
+ // expected-error at +2{{OpenACC 'num_gangs' clause cannot appear more than once on a 'kernels' directive}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels num_gangs(1) num_gangs(2)
+ while(1);
+
+ // expected-error at +2{{OpenACC 'num_gangs' clause cannot appear more than once on a 'parallel' directive}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc parallel num_gangs(1) num_gangs(2)
+ while(1);
+
+ // expected-error at +1{{too many integer expression arguments provided to OpenACC 'num_gangs' clause: 'kernels' directive expects maximum of 1, 2 were provided}}
+#pragma acc kernels num_gangs(1, getS())
+ while(1);
+
+ // expected-error at +1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}}
+#pragma acc serial num_gangs(1, getS())
+ while(1);
+#pragma acc parallel num_gangs(1, getS())
+ while(1);
+
+ struct NotConvertible{} NC;
+ // expected-error at +1{{OpenACC clause 'num_gangs' requires expression of integer type ('struct NotConvertible' invalid)}}
+#pragma acc parallel num_gangs(NC)
+ while(1);
+
+ // expected-error at +1{{OpenACC clause 'num_gangs' requires expression of integer type ('struct NotConvertible' invalid)}}
+#pragma acc parallel num_gangs(1, NC)
+ while(1);
+
+ // expected-error at +1{{OpenACC clause 'num_gangs' requires expression of integer type ('struct NotConvertible' invalid)}}
+#pragma acc parallel num_gangs(NC, 1)
+ while(1);
+
+#pragma acc parallel num_gangs(getS(), 1, getS())
+ while(1);
+
+ // expected-error at +1{{too many integer expression arguments provided to OpenACC 'num_gangs' clause: 'parallel' directive expects maximum of 3, 4 were provided}}
+#pragma acc parallel num_gangs(getS(), 1, getS(), 1)
+ while(1);
+}
diff --git a/clang/test/SemaOpenACC/compute-construct-num_gangs-clause.cpp b/clang/test/SemaOpenACC/compute-construct-num_gangs-clause.cpp
new file mode 100644
index 00000000000000..ec3df87a065572
--- /dev/null
+++ b/clang/test/SemaOpenACC/compute-construct-num_gangs-clause.cpp
@@ -0,0 +1,160 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+struct NotConvertible{} NC;
+struct Incomplete *SomeIncomplete; // #INCOMPLETE
+enum E{} SomeE;
+enum class E2{} SomeE2;
+
+struct CorrectConvert {
+ operator int();
+} Convert;
+
+struct ExplicitConvertOnly {
+ explicit operator int() const; // #EXPL_CONV
+} Explicit;
+
+struct AmbiguousConvert{
+ operator int(); // #AMBIG_INT
+ operator short(); // #AMBIG_SHORT
+ operator float();
+} Ambiguous;
+
+short some_short();
+int some_int();
+long some_long();
+
+void Test() {
+#pragma acc kernels num_gangs(1)
+ while(1);
+
+ // expected-error at +1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}}
+#pragma acc serial num_gangs(1)
+ while(1);
+
+#pragma acc parallel num_gangs(1)
+ while(1);
+
+ // expected-error at +1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}}
+#pragma acc serial num_gangs(some_short(), some_int(), some_long())
+ while(1);
+
+#pragma acc parallel num_gangs(some_short(), some_int(), some_long())
+ while(1);
+
+ // expected-error at +1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}}
+#pragma acc serial num_gangs(some_short(), some_int(), some_long(), SomeE)
+ while(1);
+
+ // expected-error at +1{{too many integer expression arguments provided to OpenACC 'num_gangs' clause: 'parallel' directive expects maximum of 3, 4 were provided}}
+#pragma acc parallel num_gangs(some_short(), some_int(), some_long(), SomeE)
+ while(1);
+
+ // expected-error at +1{{too many integer expression arguments provided to OpenACC 'num_gangs' clause: 'kernels' directive expects maximum of 1, 2 were provided}}
+#pragma acc kernels num_gangs(1, 2)
+ while(1);
+
+ // expected-error at +1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}}
+#pragma acc serial num_gangs(1, 2)
+ while(1);
+
+#pragma acc parallel num_gangs(1, 2)
+ while(1);
+
+ // expected-error at +3{{multiple conversions from expression type 'struct AmbiguousConvert' to an integral type}}
+ // expected-note@#AMBIG_INT{{conversion to integral type 'int'}}
+ // expected-note@#AMBIG_SHORT{{conversion to integral type 'short'}}
+#pragma acc parallel num_gangs(Ambiguous)
+ while(1);
+
+ // expected-error at +1{{OpenACC clause 'num_gangs' requires expression of integer type ('struct NotConvertible' invalid)}}
+#pragma acc parallel num_gangs(NC, SomeE)
+ while(1);
+
+ // expected-error at +1{{OpenACC clause 'num_gangs' requires expression of integer type ('struct NotConvertible' invalid)}}
+#pragma acc parallel num_gangs(SomeE, NC)
+ while(1);
+
+ // expected-error at +3{{OpenACC integer expression type 'struct ExplicitConvertOnly' requires explicit conversion to 'int'}}
+ // expected-note@#EXPL_CONV{{conversion to integral type 'int'}}
+ // expected-error at +1{{OpenACC clause 'num_gangs' requires expression of integer type ('struct NotConvertible' invalid)}}
+#pragma acc parallel num_gangs(Explicit, NC)
+ while(1);
+
+ // expected-error at +4{{OpenACC integer expression type 'struct ExplicitConvertOnly' requires explicit conversion to 'int'}}
+ // expected-note@#EXPL_CONV{{conversion to integral type 'int'}}
+ // expected-error at +2{{OpenACC clause 'num_gangs' requires expression of integer type ('struct NotConvertible' invalid)}}
+ // expected-error at +1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}}
+#pragma acc serial num_gangs(Explicit, NC)
+ while(1);
+
+ // expected-error at +6{{OpenACC integer expression type 'struct ExplicitConvertOnly' requires explicit conversion to 'int'}}
+ // expected-note@#EXPL_CONV{{conversion to integral type 'int'}}
+ // expected-error at +4{{OpenACC clause 'num_gangs' requires expression of integer type ('struct NotConvertible' invalid)}}
+ // expected-error at +3{{multiple conversions from expression type 'struct AmbiguousConvert' to an integral type}}
+ // expected-note@#AMBIG_INT{{conversion to integral type 'int'}}
+ // expected-note@#AMBIG_SHORT{{conversion to integral type 'short'}}
+#pragma acc parallel num_gangs(Explicit, NC, Ambiguous)
+ while(1);
+
+ // expected-error at +7{{OpenACC integer expression type 'struct ExplicitConvertOnly' requires explicit conversion to 'int'}}
+ // expected-note@#EXPL_CONV{{conversion to integral type 'int'}}
+ // expected-error at +5{{OpenACC clause 'num_gangs' requires expression of integer type ('struct NotConvertible' invalid)}}
+ // expected-error at +4{{multiple conversions from expression type 'struct AmbiguousConvert' to an integral type}}
+ // expected-note@#AMBIG_INT{{conversion to integral type 'int'}}
+ // expected-note@#AMBIG_SHORT{{conversion to integral type 'short'}}
+ // expected-error at +1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}}
+#pragma acc serial num_gangs(Explicit, NC, Ambiguous)
+ while(1);
+ // TODO
+}
+
+struct HasInt {
+ using IntTy = int;
+ using ShortTy = short;
+ static constexpr int value = 1;
+ static constexpr AmbiguousConvert ACValue;
+ static constexpr ExplicitConvertOnly EXValue;
+
+ operator char();
+};
+
+template <typename T>
+void TestInst() {
+ // expected-error at +2{{no member named 'Invalid' in 'HasInt'}}
+ // expected-error at +1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}}
+#pragma acc serial num_gangs(HasInt::Invalid)
+ while(1);
+
+ // expected-error at +2{{no member named 'Invalid' in 'HasInt'}}
+ // expected-note@#INST{{in instantiation of function template specialization}}
+#pragma acc parallel num_gangs(T::Invalid)
+ while(1);
+
+ // expected-error at +1{{no member named 'Invalid' in 'HasInt'}}
+#pragma acc parallel num_gangs(1, HasInt::Invalid)
+ while(1);
+
+ // expected-error at +1{{no member named 'Invalid' in 'HasInt'}}
+#pragma acc parallel num_gangs(T::Invalid, 1)
+ while(1);
+
+ // expected-error at +2{{no member named 'Invalid' in 'HasInt'}}
+ // expected-error at +1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}}
+#pragma acc serial num_gangs(1, HasInt::Invalid)
+ while(1);
+
+ // expected-error at +1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}}
+#pragma acc serial num_gangs(T::Invalid, 1)
+ while(1);
+
+#pragma acc parallel num_gangs(T::value, typename T::IntTy{})
+ while(1);
+
+ // expected-error at +1{{OpenACC 'num_gangs' clause is not valid on 'serial' directive}}
+#pragma acc serial num_gangs(T::value, typename T::IntTy{})
+ while(1);
+}
+
+void Inst() {
+ TestInst<HasInt>(); // #INST
+}
diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
index cbc1d85bb33dfc..74163f30e19b1d 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -2803,6 +2803,10 @@ void OpenACCClauseEnqueue::VisitVectorLengthClause(
const OpenACCVectorLengthClause &C) {
Visitor.AddStmt(C.getIntExpr());
}
+void OpenACCClauseEnqueue::VisitNumGangsClause(const OpenACCNumGangsClause &C) {
+ for (Expr *IE : C.getIntExprs())
+ Visitor.AddStmt(IE);
+}
} // namespace
void EnqueueVisitor::EnqueueChildren(const OpenACCClause *C) {
More information about the cfe-commits
mailing list