[clang] [OpenACC] Implement `self` clause for compute constructs (PR #88760)
Erich Keane via cfe-commits
cfe-commits at lists.llvm.org
Mon Apr 15 10:28:14 PDT 2024
https://github.com/erichkeane created https://github.com/llvm/llvm-project/pull/88760
`self` clauses on compute constructs take an optional condition expression. We again limit the implementation to ONLY compute constructs to ensure we get all the rules correct for others. However, this one will be particularly complicated, as it takes a `var-list` for `update`, so when we get to that construct/clause combination, we need to do that as well.
This patch also furthers uses of the `OpenACCClauses.def` as it became useful while implementing this (as well as some other minor refactors as I went through).
Finally, `self` and `if` clauses have an interaction with each other, if an `if` clause evaluates to `true`, the `self` clause has no effect. While this is intended and can be used 'meaningfully', we are warning on this with a very granular warning, so that this edge case will be noticed by newer users, but can be disabled trivially.
>From a96f008c7e4d97ca634f8ce24fc43843d5e19c91 Mon Sep 17 00:00:00 2001
From: erichkeane <ekeane at nvidia.com>
Date: Fri, 12 Apr 2024 10:49:17 -0700
Subject: [PATCH] [OpenACC] Implement `self` clause for compute constructs
`self` clauses on compute constructs take an optional condition
expression. We again limit the implementation to ONLY compute constructs
to ensure we get all the rules correct for others. However, this one
will be particularly complicated, as it takes a `var-list` for `update`,
so when we get to that construct/clause combination, we need to do that
as well.
This patch also furthers uses of the `OpenACCClauses.def` as it became
useful while implementing this (as well as some other minor refactors as
I went through).
Finally, `self` and `if` clauses have an interaction with each other, if
an `if` clause evaluates to `true`, the `self` clause has no effect.
While this is intended and can be used 'meaningfully', we are warning on
this with a very granular warning, so that this edge case will be
noticed by newer users, but can be disabled trivially.
---
clang/include/clang/AST/OpenACCClause.h | 65 ++++--------
clang/include/clang/AST/StmtOpenACC.h | 4 +-
.../clang/Basic/DiagnosticSemaKinds.td | 4 +
clang/include/clang/Basic/OpenACCClauses.def | 1 +
clang/include/clang/Basic/OpenACCKinds.h | 6 ++
clang/include/clang/Sema/SemaOpenACC.h | 18 +++-
clang/lib/AST/OpenACCClause.cpp | 26 +++++
clang/lib/AST/StmtProfile.cpp | 5 +
clang/lib/AST/TextNodeDumper.cpp | 1 +
clang/lib/Parse/ParseOpenACC.cpp | 15 ++-
clang/lib/Sema/SemaOpenACC.cpp | 71 +++++++++++--
clang/lib/Sema/TreeTransform.h | 100 +++++++++++++-----
clang/lib/Serialization/ASTReader.cpp | 7 +-
clang/lib/Serialization/ASTWriter.cpp | 9 +-
clang/test/ParserOpenACC/parse-clauses.c | 5 +-
.../compute-construct-clause-ast.cpp | 91 ++++++++++++++++
.../compute-construct-self-clause.c | 82 ++++++++++++++
.../compute-construct-self-clause.cpp | 99 +++++++++++++++++
clang/tools/libclang/CIndex.cpp | 4 +
19 files changed, 516 insertions(+), 97 deletions(-)
create mode 100644 clang/test/SemaOpenACC/compute-construct-self-clause.c
create mode 100644 clang/test/SemaOpenACC/compute-construct-self-clause.cpp
diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h
index 401b8e904a1b7a..07587849eb1219 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -145,6 +145,17 @@ class OpenACCIfClause : public OpenACCClauseWithCondition {
SourceLocation EndLoc);
};
+/// A 'self' clause, which has an optional condition expression.
+class OpenACCSelfClause : public OpenACCClauseWithCondition {
+ OpenACCSelfClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
+ Expr *ConditionExpr, SourceLocation EndLoc);
+
+public:
+ static OpenACCSelfClause *Create(const ASTContext &C, SourceLocation BeginLoc,
+ SourceLocation LParenLoc,
+ Expr *ConditionExpr, SourceLocation EndLoc);
+};
+
template <class Impl> class OpenACCClauseVisitor {
Impl &getDerived() { return static_cast<Impl &>(*this); }
@@ -159,53 +170,13 @@ template <class Impl> class OpenACCClauseVisitor {
return;
switch (C->getClauseKind()) {
- case OpenACCClauseKind::Default:
- VisitDefaultClause(*cast<OpenACCDefaultClause>(C));
- return;
- case OpenACCClauseKind::If:
- VisitIfClause(*cast<OpenACCIfClause>(C));
- return;
- case OpenACCClauseKind::Finalize:
- case OpenACCClauseKind::IfPresent:
- case OpenACCClauseKind::Seq:
- case OpenACCClauseKind::Independent:
- case OpenACCClauseKind::Auto:
- case OpenACCClauseKind::Worker:
- case OpenACCClauseKind::Vector:
- case OpenACCClauseKind::NoHost:
- case OpenACCClauseKind::Self:
- case OpenACCClauseKind::Copy:
- case OpenACCClauseKind::UseDevice:
- case OpenACCClauseKind::Attach:
- case OpenACCClauseKind::Delete:
- case OpenACCClauseKind::Detach:
- case OpenACCClauseKind::Device:
- case OpenACCClauseKind::DevicePtr:
- case OpenACCClauseKind::DeviceResident:
- case OpenACCClauseKind::FirstPrivate:
- case OpenACCClauseKind::Host:
- case OpenACCClauseKind::Link:
- case OpenACCClauseKind::NoCreate:
- case OpenACCClauseKind::Present:
- case OpenACCClauseKind::Private:
- case OpenACCClauseKind::CopyOut:
- case OpenACCClauseKind::CopyIn:
- case OpenACCClauseKind::Create:
- case OpenACCClauseKind::Reduction:
- case OpenACCClauseKind::Collapse:
- case OpenACCClauseKind::Bind:
- case OpenACCClauseKind::VectorLength:
- case OpenACCClauseKind::NumGangs:
- case OpenACCClauseKind::NumWorkers:
- case OpenACCClauseKind::DeviceNum:
- case OpenACCClauseKind::DefaultAsync:
- case OpenACCClauseKind::DeviceType:
- case OpenACCClauseKind::DType:
- case OpenACCClauseKind::Async:
- case OpenACCClauseKind::Tile:
- case OpenACCClauseKind::Gang:
- case OpenACCClauseKind::Wait:
- case OpenACCClauseKind::Invalid:
+#define VISIT_CLAUSE(CLAUSE_NAME) \
+ case OpenACCClauseKind::CLAUSE_NAME: \
+ Visit##CLAUSE_NAME##Clause(*cast<OpenACC##CLAUSE_NAME##Clause>(C)); \
+ return;
+#include "clang/Basic/OpenACCClauses.def"
+
+ default:
llvm_unreachable("Clause visitor not yet implemented");
}
llvm_unreachable("Invalid Clause kind");
diff --git a/clang/include/clang/AST/StmtOpenACC.h b/clang/include/clang/AST/StmtOpenACC.h
index 419cb6cada0bc7..66f8f844e0b29e 100644
--- a/clang/include/clang/AST/StmtOpenACC.h
+++ b/clang/include/clang/AST/StmtOpenACC.h
@@ -142,9 +142,7 @@ class OpenACCComputeConstruct final
Stmt *StructuredBlock)
: OpenACCAssociatedStmtConstruct(OpenACCComputeConstructClass, K, Start,
End, StructuredBlock) {
- assert((K == OpenACCDirectiveKind::Parallel ||
- K == OpenACCDirectiveKind::Serial ||
- K == OpenACCDirectiveKind::Kernels) &&
+ assert(isOpenACCComputeDirectiveKind(K) &&
"Only parallel, serial, and kernels constructs should be "
"represented by this type");
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 5ec0218aedfe86..44f802c0c28e84 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12274,4 +12274,8 @@ def note_acc_branch_into_compute_construct
: Note<"invalid branch into OpenACC Compute Construct">;
def note_acc_branch_out_of_compute_construct
: Note<"invalid branch out of OpenACC Compute Construct">;
+def warn_acc_if_self_conflict
+ : Warning<"OpenACC construct 'self' has no effect when an 'if' clause "
+ "evaluates to true">,
+ InGroup<DiagGroup<"openacc-self-if-potential-conflict">>;
} // end of sema component.
diff --git a/clang/include/clang/Basic/OpenACCClauses.def b/clang/include/clang/Basic/OpenACCClauses.def
index 7fd2720e02ce22..378495d2c0909a 100644
--- a/clang/include/clang/Basic/OpenACCClauses.def
+++ b/clang/include/clang/Basic/OpenACCClauses.def
@@ -17,5 +17,6 @@
VISIT_CLAUSE(Default)
VISIT_CLAUSE(If)
+VISIT_CLAUSE(Self)
#undef VISIT_CLAUSE
diff --git a/clang/include/clang/Basic/OpenACCKinds.h b/clang/include/clang/Basic/OpenACCKinds.h
index 3414df99991701..e3f74178433285 100644
--- a/clang/include/clang/Basic/OpenACCKinds.h
+++ b/clang/include/clang/Basic/OpenACCKinds.h
@@ -146,6 +146,12 @@ inline llvm::raw_ostream &operator<<(llvm::raw_ostream &Out,
return printOpenACCDirectiveKind(Out, K);
}
+inline bool isOpenACCComputeDirectiveKind(OpenACCDirectiveKind K) {
+ return K == OpenACCDirectiveKind::Parallel ||
+ K == OpenACCDirectiveKind::Serial ||
+ K == OpenACCDirectiveKind::Kernels;
+}
+
enum class OpenACCAtomicKind {
Read,
Write,
diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h
index c1fe0f5b9c0f6b..329dc3945fa2a6 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -44,7 +44,8 @@ class SemaOpenACC : public SemaBase {
Expr *ConditionExpr;
};
- std::variant<DefaultDetails, ConditionDetails> Details;
+ std::variant<std::monostate, DefaultDetails, ConditionDetails> Details =
+ std::monostate{};
public:
OpenACCParsedClause(OpenACCDirectiveKind DirKind,
@@ -72,8 +73,17 @@ class SemaOpenACC : public SemaBase {
}
Expr *getConditionExpr() {
- assert(ClauseKind == OpenACCClauseKind::If &&
+ assert((ClauseKind == OpenACCClauseKind::If ||
+ (ClauseKind == OpenACCClauseKind::Self &&
+ DirKind != OpenACCDirectiveKind::Update)) &&
"Parsed clause kind does not have a condition expr");
+
+ // 'self' has an optional ConditionExpr, so be tolerant of that. This will
+ // assert in variant otherwise.
+ if (ClauseKind == OpenACCClauseKind::Self &&
+ std::holds_alternative<std::monostate>(Details))
+ return nullptr;
+
return std::get<ConditionDetails>(Details).ConditionExpr;
}
@@ -87,7 +97,9 @@ class SemaOpenACC : public SemaBase {
}
void setConditionDetails(Expr *ConditionExpr) {
- assert(ClauseKind == OpenACCClauseKind::If &&
+ assert((ClauseKind == OpenACCClauseKind::If ||
+ (ClauseKind == OpenACCClauseKind::Self &&
+ DirKind != OpenACCDirectiveKind::Update)) &&
"Parsed clause kind does not have a condition expr");
// In C++ we can count on this being a 'bool', but in C this gets left as
// some sort of scalar that codegen will have to take care of converting.
diff --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp
index dcb512cb514179..9c259c8f9bd0a1 100644
--- a/clang/lib/AST/OpenACCClause.cpp
+++ b/clang/lib/AST/OpenACCClause.cpp
@@ -48,6 +48,26 @@ OpenACCIfClause::OpenACCIfClause(SourceLocation BeginLoc,
"Condition expression type not scalar/dependent");
}
+OpenACCSelfClause *OpenACCSelfClause::Create(const ASTContext &C,
+ SourceLocation BeginLoc,
+ SourceLocation LParenLoc,
+ Expr *ConditionExpr,
+ SourceLocation EndLoc) {
+ void *Mem = C.Allocate(sizeof(OpenACCIfClause), alignof(OpenACCIfClause));
+ return new (Mem)
+ OpenACCSelfClause(BeginLoc, LParenLoc, ConditionExpr, EndLoc);
+}
+
+OpenACCSelfClause::OpenACCSelfClause(SourceLocation BeginLoc,
+ SourceLocation LParenLoc,
+ Expr *ConditionExpr, SourceLocation EndLoc)
+ : OpenACCClauseWithCondition(OpenACCClauseKind::Self, BeginLoc, LParenLoc,
+ ConditionExpr, EndLoc) {
+ assert((!ConditionExpr || ConditionExpr->isInstantiationDependent() ||
+ ConditionExpr->getType()->isScalarType()) &&
+ "Condition expression type not scalar/dependent");
+}
+
OpenACCClause::child_range OpenACCClause::children() {
switch (getClauseKind()) {
default:
@@ -72,3 +92,9 @@ void OpenACCClausePrinter::VisitDefaultClause(const OpenACCDefaultClause &C) {
void OpenACCClausePrinter::VisitIfClause(const OpenACCIfClause &C) {
OS << "if(" << C.getConditionExpr() << ")";
}
+
+void OpenACCClausePrinter::VisitSelfClause(const OpenACCSelfClause &C) {
+ OS << "self";
+ if (const Expr *CondExpr = C.getConditionExpr())
+ OS << "(" << CondExpr << ")";
+}
diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index d2aac1e640380f..6b9292e1d0da39 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -2473,6 +2473,11 @@ void OpenACCClauseProfiler::VisitIfClause(const OpenACCIfClause &Clause) {
"if clause requires a valid condition expr");
Profiler.VisitStmt(Clause.getConditionExpr());
}
+
+void OpenACCClauseProfiler::VisitSelfClause(const OpenACCSelfClause &Clause) {
+ if (Clause.hasConditionExpr())
+ Profiler.VisitStmt(Clause.getConditionExpr());
+}
} // namespace
void StmtProfiler::VisitOpenACCComputeConstruct(
diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp
index 688daa64d61974..ff5b3df2d6dfac 100644
--- a/clang/lib/AST/TextNodeDumper.cpp
+++ b/clang/lib/AST/TextNodeDumper.cpp
@@ -398,6 +398,7 @@ void TextNodeDumper::Visit(const OpenACCClause *C) {
OS << '(' << cast<OpenACCDefaultClause>(C)->getDefaultClauseKind() << ')';
break;
case OpenACCClauseKind::If:
+ case OpenACCClauseKind::Self:
// The condition expression will be printed as a part of the 'children',
// but print 'clause' here so it is clear what is happening from the dump.
OS << " clause";
diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp
index 91f2b8afcf0c24..da7132b5c4226f 100644
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@@ -835,18 +835,23 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
case OpenACCClauseKind::Default: {
Token DefKindTok = getCurToken();
- if (expectIdentifierOrKeyword(*this))
- break;
+ if (expectIdentifierOrKeyword(*this)) {
+ Parens.skipToEnd();
+ return OpenACCCanContinue();
+ }
ConsumeToken();
OpenACCDefaultClauseKind DefKind =
getOpenACCDefaultClauseKind(DefKindTok);
- if (DefKind == OpenACCDefaultClauseKind::Invalid)
+ if (DefKind == OpenACCDefaultClauseKind::Invalid) {
Diag(DefKindTok, diag::err_acc_invalid_default_clause_kind);
- else
+ Parens.skipToEnd();
+ return OpenACCCanContinue();
+ } else {
ParsedClause.setDefaultDetails(DefKind);
+ }
break;
}
@@ -977,6 +982,8 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
case OpenACCClauseKind::Self: {
assert(DirKind != OpenACCDirectiveKind::Update);
ExprResult CondExpr = ParseOpenACCConditionExpr();
+ ParsedClause.setConditionDetails(CondExpr.isUsable() ? CondExpr.get()
+ : nullptr);
if (CondExpr.isInvalid()) {
Parens.skipToEnd();
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 1249136c87650b..04d9925327a8a3 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -76,6 +76,19 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
default:
return false;
}
+ case OpenACCClauseKind::Self:
+ switch (DirectiveKind) {
+ case OpenACCDirectiveKind::Parallel:
+ case OpenACCDirectiveKind::Serial:
+ case OpenACCDirectiveKind::Kernels:
+ case OpenACCDirectiveKind::Update:
+ case OpenACCDirectiveKind::ParallelLoop:
+ case OpenACCDirectiveKind::SerialLoop:
+ case OpenACCDirectiveKind::KernelsLoop:
+ return true;
+ default:
+ return false;
+ }
default:
// Do nothing so we can go to the 'unimplemented' diagnostic instead.
return true;
@@ -121,9 +134,7 @@ SemaOpenACC::ActOnClause(ArrayRef<const OpenACCClause *> ExistingClauses,
// 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 (Clause.getDirectiveKind() != OpenACCDirectiveKind::Parallel &&
- Clause.getDirectiveKind() != OpenACCDirectiveKind::Serial &&
- Clause.getDirectiveKind() != OpenACCDirectiveKind::Kernels)
+ if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()))
break;
// Don't add an invalid clause to the AST.
@@ -146,9 +157,7 @@ SemaOpenACC::ActOnClause(ArrayRef<const OpenACCClause *> ExistingClauses,
// 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 (Clause.getDirectiveKind() != OpenACCDirectiveKind::Parallel &&
- Clause.getDirectiveKind() != OpenACCDirectiveKind::Serial &&
- Clause.getDirectiveKind() != OpenACCDirectiveKind::Kernels)
+ if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()))
break;
// There is no prose in the standard that says duplicates aren't allowed,
@@ -160,12 +169,58 @@ SemaOpenACC::ActOnClause(ArrayRef<const OpenACCClause *> ExistingClauses,
// The parser has ensured that we have a proper condition expr, so there
// isn't really much to do here.
- // TODO OpenACC: When we implement 'self', this clauses causes us to
- // 'ignore' the self clause, so we should implement a warning here.
+ // If the 'if' clause is true, it makes the 'self' clause have no effect,
+ // diagnose that here.
+ // TODO OpenACC: When we add these two to other constructs, we might not
+ // want to warn on this (for example, 'update').
+ const auto *Itr =
+ llvm::find_if(ExistingClauses, [](const OpenACCClause *C) {
+ return C->getClauseKind() == OpenACCClauseKind::Self;
+ });
+ if (Itr != ExistingClauses.end()) {
+ Diag(Clause.getBeginLoc(), diag::warn_acc_if_self_conflict);
+ Diag((*Itr)->getBeginLoc(), diag::note_acc_previous_clause_here);
+ }
+
return OpenACCIfClause::Create(
getASTContext(), Clause.getBeginLoc(), Clause.getLParenLoc(),
Clause.getConditionExpr(), Clause.getEndLoc());
}
+
+ case OpenACCClauseKind::Self: {
+ // 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;
+
+ // TODO OpenACC: When we implement this for 'update', this takes a
+ // 'var-list' instead of a condition expression, so semantics/handling has
+ // to happen differently here.
+
+ // 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 the 'if' clause is true, it makes the 'self' clause have no effect,
+ // diagnose that here.
+ // TODO OpenACC: When we add these two to other constructs, we might not
+ // want to warn on this (for example, 'update').
+ const auto *Itr =
+ llvm::find_if(ExistingClauses, [](const OpenACCClause *C) {
+ return C->getClauseKind() == OpenACCClauseKind::If;
+ });
+ if (Itr != ExistingClauses.end()) {
+ Diag(Clause.getBeginLoc(), diag::warn_acc_if_self_conflict);
+ Diag((*Itr)->getBeginLoc(), diag::note_acc_previous_clause_here);
+ }
+
+ return OpenACCSelfClause::Create(
+ getASTContext(), Clause.getBeginLoc(), Clause.getLParenLoc(),
+ Clause.getConditionExpr(), Clause.getEndLoc());
+ }
default:
break;
}
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index fdf84c512c4931..d2ef6b153e4ab3 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -11080,6 +11080,77 @@ OMPClause *TreeTransform<Derived>::TransformOMPXBareClause(OMPXBareClause *C) {
//===----------------------------------------------------------------------===//
// OpenACC transformation
//===----------------------------------------------------------------------===//
+namespace {
+template <typename Derived>
+class OpenACCClauseTransform final
+ : public OpenACCClauseVisitor<OpenACCClauseTransform<Derived>> {
+ TreeTransform<Derived> &Self;
+ SemaOpenACC::OpenACCParsedClause &ParsedClause;
+ OpenACCClause *NewClause = nullptr;
+
+public:
+ OpenACCClauseTransform(TreeTransform<Derived> &Self,
+ SemaOpenACC::OpenACCParsedClause &PC)
+ : Self(Self), ParsedClause(PC) {}
+
+ OpenACCClause *CreatedClause() const { return NewClause; }
+
+#define VISIT_CLAUSE(CLAUSE_NAME) \
+ void Visit##CLAUSE_NAME##Clause(const OpenACC##CLAUSE_NAME##Clause &Clause);
+#include "clang/Basic/OpenACCClauses.def"
+};
+
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitDefaultClause(
+ const OpenACCDefaultClause &C) {
+ ParsedClause.setDefaultDetails(C.getDefaultClauseKind());
+
+ NewClause = OpenACCDefaultClause::Create(
+ Self.getSema().getASTContext(), ParsedClause.getDefaultClauseKind(),
+ ParsedClause.getBeginLoc(), ParsedClause.getLParenLoc(),
+ ParsedClause.getEndLoc());
+}
+
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitIfClause(const OpenACCIfClause &C) {
+ Expr *Cond = const_cast<Expr *>(C.getConditionExpr());
+ assert(Cond && "If constructed with invalid Condition");
+ Sema::ConditionResult Res = Self.TransformCondition(
+ Cond->getExprLoc(), /*Var=*/nullptr, Cond, Sema::ConditionKind::Boolean);
+
+ if (Res.isInvalid() || !Res.get().second)
+ return;
+
+ ParsedClause.setConditionDetails(Res.get().second);
+
+ NewClause = OpenACCIfClause::Create(
+ Self.getSema().getASTContext(), ParsedClause.getBeginLoc(),
+ ParsedClause.getLParenLoc(), ParsedClause.getConditionExpr(),
+ ParsedClause.getEndLoc());
+}
+
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitSelfClause(
+ const OpenACCSelfClause &C) {
+
+ if (C.hasConditionExpr()) {
+ Expr *Cond = const_cast<Expr *>(C.getConditionExpr());
+ Sema::ConditionResult Res =
+ Self.TransformCondition(Cond->getExprLoc(), /*Var=*/nullptr, Cond,
+ Sema::ConditionKind::Boolean);
+
+ if (Res.isInvalid() || !Res.get().second)
+ return;
+
+ ParsedClause.setConditionDetails(Res.get().second);
+ }
+
+ NewClause = OpenACCSelfClause::Create(
+ Self.getSema().getASTContext(), ParsedClause.getBeginLoc(),
+ ParsedClause.getLParenLoc(), ParsedClause.getConditionExpr(),
+ ParsedClause.getEndLoc());
+}
+} // namespace
template <typename Derived>
OpenACCClause *TreeTransform<Derived>::TransformOpenACCClause(
ArrayRef<const OpenACCClause *> ExistingClauses,
@@ -11092,33 +11163,10 @@ OpenACCClause *TreeTransform<Derived>::TransformOpenACCClause(
if (const auto *WithParms = dyn_cast<OpenACCClauseWithParams>(OldClause))
ParsedClause.setLParenLoc(WithParms->getLParenLoc());
- switch (OldClause->getClauseKind()) {
- case OpenACCClauseKind::Default:
- // There is nothing to do here as nothing dependent can appear in this
- // clause. So just set the values so Sema can set the right value.
- ParsedClause.setDefaultDetails(
- cast<OpenACCDefaultClause>(OldClause)->getDefaultClauseKind());
- break;
- case OpenACCClauseKind::If: {
- Expr *Cond = const_cast<Expr *>(
- cast<OpenACCIfClause>(OldClause)->getConditionExpr());
- assert(Cond && "If constructed with invalid Condition");
- Sema::ConditionResult Res =
- TransformCondition(Cond->getExprLoc(), /*Var=*/nullptr, Cond,
- Sema::ConditionKind::Boolean);
-
- if (Res.isInvalid() || !Res.get().second)
- return nullptr;
-
- ParsedClause.setConditionDetails(Res.get().second);
- break;
- }
- default:
- assert(false && "Unhandled OpenACC clause in TreeTransform");
- return nullptr;
- }
+ OpenACCClauseTransform<Derived> Transform{*this, ParsedClause};
+ Transform.Visit(OldClause);
- return getSema().OpenACC().ActOnClause(ExistingClauses, ParsedClause);
+ return Transform.CreatedClause();
}
template <typename Derived>
diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
index 8c4b460970ad2b..6f404435896159 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -11794,6 +11794,12 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
return OpenACCIfClause::Create(getContext(), BeginLoc, LParenLoc, CondExpr,
EndLoc);
}
+ case OpenACCClauseKind::Self: {
+ SourceLocation LParenLoc = readSourceLocation();
+ Expr *CondExpr = readBool() ? readSubExpr() : nullptr;
+ return OpenACCSelfClause::Create(getContext(), BeginLoc, LParenLoc,
+ CondExpr, EndLoc);
+ }
case OpenACCClauseKind::Finalize:
case OpenACCClauseKind::IfPresent:
case OpenACCClauseKind::Seq:
@@ -11802,7 +11808,6 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
case OpenACCClauseKind::Worker:
case OpenACCClauseKind::Vector:
case OpenACCClauseKind::NoHost:
- case OpenACCClauseKind::Self:
case OpenACCClauseKind::Copy:
case OpenACCClauseKind::UseDevice:
case OpenACCClauseKind::Attach:
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index 85b7fd5535a1bf..0459145860135a 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -7524,6 +7524,14 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
AddStmt(const_cast<Expr *>(IC->getConditionExpr()));
return;
}
+ case OpenACCClauseKind::Self: {
+ const auto *SC = cast<OpenACCIfClause>(C);
+ writeSourceLocation(SC->getLParenLoc());
+ writeBool(SC->hasConditionExpr());
+ if (SC->hasConditionExpr())
+ AddStmt(const_cast<Expr *>(SC->getConditionExpr()));
+ return;
+ }
case OpenACCClauseKind::Finalize:
case OpenACCClauseKind::IfPresent:
case OpenACCClauseKind::Seq:
@@ -7532,7 +7540,6 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
case OpenACCClauseKind::Worker:
case OpenACCClauseKind::Vector:
case OpenACCClauseKind::NoHost:
- case OpenACCClauseKind::Self:
case OpenACCClauseKind::Copy:
case OpenACCClauseKind::UseDevice:
case OpenACCClauseKind::Attach:
diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c
index 2369df58308a72..4462f0df540f2d 100644
--- a/clang/test/ParserOpenACC/parse-clauses.c
+++ b/clang/test/ParserOpenACC/parse-clauses.c
@@ -376,16 +376,13 @@ void SelfClause() {
#pragma acc serial self(i > j, seq
for(;;){}
- // expected-warning at +2{{left operand of comma operator has no effect}}
- // expected-warning at +1{{OpenACC clause 'self' not yet implemented, clause ignored}}
+ // expected-warning at +1{{left operand of comma operator has no effect}}
#pragma acc serial self(i, j)
for(;;){}
- // expected-warning at +1{{OpenACC clause 'self' not yet implemented, clause ignored}}
#pragma acc serial self(i > j)
for(;;){}
- // expected-warning at +2{{OpenACC clause 'self' not yet implemented, clause ignored}}
// expected-warning at +1{{OpenACC clause 'seq' not yet implemented, clause ignored}}
#pragma acc serial self(1+5>3), seq
for(;;){}
diff --git a/clang/test/SemaOpenACC/compute-construct-clause-ast.cpp b/clang/test/SemaOpenACC/compute-construct-clause-ast.cpp
index 018f0b68c78109..6d2efcf81eb6e4 100644
--- a/clang/test/SemaOpenACC/compute-construct-clause-ast.cpp
+++ b/clang/test/SemaOpenACC/compute-construct-clause-ast.cpp
@@ -110,6 +110,50 @@ void TemplFunc() {
// CHECK-NEXT: CXXBoolLiteralExpr
// CHECK-NEXT: NullStmt
+#pragma acc serial self
+ while(true);
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}}serial
+ // CHECK-NEXT: self clause
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: NullStmt
+
+#pragma acc kernels self(T::SomeFloat)
+ while(true);
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels
+ // CHECK-NEXT: self clause
+ // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue
+ // CHECK-NEXT: NestedNameSpecifier TypeSpec 'T'
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: NullStmt
+
+#pragma acc parallel self(T::SomeFloat) if (T::SomeFloat)
+ while(true);
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
+ // CHECK-NEXT: self clause
+ // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue
+ // CHECK-NEXT: NestedNameSpecifier TypeSpec 'T'
+ // CHECK-NEXT: if clause
+ // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue
+ // CHECK-NEXT: NestedNameSpecifier TypeSpec 'T'
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: NullStmt
+
+#pragma acc serial if(T::SomeFloat) self(T::SomeFloat)
+ while(true);
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}}serial
+ // CHECK-NEXT: if clause
+ // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue
+ // CHECK-NEXT: NestedNameSpecifier TypeSpec 'T'
+ // CHECK-NEXT: self clause
+ // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue
+ // CHECK-NEXT: NestedNameSpecifier TypeSpec 'T'
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: NullStmt
+
// Match the instantiation:
// CHECK: FunctionDecl{{.*}}TemplFunc{{.*}}implicit_instantiation
// CHECK-NEXT: TemplateArgument type 'InstTy'
@@ -171,6 +215,53 @@ void TemplFunc() {
// CHECK-NEXT: WhileStmt
// CHECK-NEXT: CXXBoolLiteralExpr
// CHECK-NEXT: NullStmt
+
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}}serial
+ // CHECK-NEXT: self clause
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: NullStmt
+
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels
+ // CHECK-NEXT: self clause
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'bool' <FloatingToBoolean>
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'float' <LValueToRValue>
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'const float' lvalue Var{{.*}} 'SomeFloat' 'const float'
+ // CHECK-NEXT: NestedNameSpecifier TypeSpec 'InstTy'
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: NullStmt
+
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
+ // CHECK-NEXT: self clause
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'bool' <FloatingToBoolean>
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'float' <LValueToRValue>
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'const float' lvalue Var{{.*}} 'SomeFloat' 'const float'
+ // CHECK-NEXT: NestedNameSpecifier TypeSpec 'InstTy'
+ // CHECK-NEXT: if clause
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'bool' <FloatingToBoolean>
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'float' <LValueToRValue>
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'const float' lvalue Var{{.*}} 'SomeFloat' 'const float'
+ // CHECK-NEXT: NestedNameSpecifier TypeSpec 'InstTy'
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: NullStmt
+
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}}serial
+ // CHECK-NEXT: if clause
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'bool' <FloatingToBoolean>
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'float' <LValueToRValue>
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'const float' lvalue Var{{.*}} 'SomeFloat' 'const float'
+ // CHECK-NEXT: NestedNameSpecifier TypeSpec 'InstTy'
+ // CHECK-NEXT: self clause
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'bool' <FloatingToBoolean>
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'float' <LValueToRValue>
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'const float' lvalue Var{{.*}} 'SomeFloat' 'const float'
+ // CHECK-NEXT: NestedNameSpecifier TypeSpec 'InstTy'
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: NullStmt
+
}
struct BoolConversion{ operator bool() const;};
diff --git a/clang/test/SemaOpenACC/compute-construct-self-clause.c b/clang/test/SemaOpenACC/compute-construct-self-clause.c
new file mode 100644
index 00000000000000..fbed2953419a2e
--- /dev/null
+++ b/clang/test/SemaOpenACC/compute-construct-self-clause.c
@@ -0,0 +1,82 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+void BoolExpr(int *I, float *F) {
+ typedef struct {} SomeStruct;
+ struct C{};
+ // expected-error at +1{{expected expression}}
+#pragma acc parallel self (struct C f())
+ while(0);
+
+ // expected-error at +1{{unexpected type name 'SomeStruct': expected expression}}
+#pragma acc serial self (SomeStruct)
+ while(0);
+
+ // expected-error at +1{{unexpected type name 'SomeStruct': expected expression}}
+#pragma acc serial self (SomeStruct())
+ while(0);
+
+ SomeStruct S;
+ // expected-error at +1{{statement requires expression of scalar type ('SomeStruct' invalid)}}
+#pragma acc serial self (S)
+ while(0);
+
+#pragma acc parallel self (I)
+ while(0);
+
+#pragma acc serial self (F)
+ while(0);
+
+#pragma acc kernels self (*I < *F)
+ while(0);
+}
+
+void WarnMaybeNotUsed(int val1, int val2) {
+
+ // expected-warning at +2{{OpenACC construct 'self' has no effect when an 'if' clause evaluates to true}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc parallel self if(val1)
+ while(0);
+
+ // expected-warning at +2{{OpenACC construct 'self' has no effect when an 'if' clause evaluates to true}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc parallel self(val1) if(val1)
+ while(0);
+
+ // expected-warning at +2{{OpenACC construct 'self' has no effect when an 'if' clause evaluates to true}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc parallel if(val1) self
+ while(0);
+
+ // expected-warning at +2{{OpenACC construct 'self' has no effect when an 'if' clause evaluates to true}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc parallel if(val1) self(val2)
+ while(0);
+
+ // The below don't warn because one side or the other has an error, thus is
+ // not added to the AST.
+
+ // expected-error at +1{{use of undeclared identifier 'invalid'}}
+#pragma acc parallel self if(invalid)
+ while(0);
+
+ // expected-error at +1{{use of undeclared identifier 'invalid'}}
+#pragma acc parallel self(invalid) if(val1)
+ while(0);
+
+ // expected-error at +2{{expected expression}}
+ // expected-error at +1{{use of undeclared identifier 'invalid'}}
+#pragma acc parallel self() if(invalid)
+ while(0);
+
+ // expected-error at +1{{use of undeclared identifier 'invalid'}}
+#pragma acc parallel if(invalid) self
+ while(0);
+
+ // expected-error at +1{{use of undeclared identifier 'invalid'}}
+#pragma acc parallel if(val2) self(invalid)
+ while(0);
+
+ // expected-error at +1{{use of undeclared identifier 'invalid'}}
+#pragma acc parallel if(invalid) self(val1)
+ while(0);
+}
diff --git a/clang/test/SemaOpenACC/compute-construct-self-clause.cpp b/clang/test/SemaOpenACC/compute-construct-self-clause.cpp
new file mode 100644
index 00000000000000..60edbdc2b1191b
--- /dev/null
+++ b/clang/test/SemaOpenACC/compute-construct-self-clause.cpp
@@ -0,0 +1,99 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+struct NoBoolConversion{};
+struct BoolConversion{
+ operator bool();
+};
+
+template <typename T, typename U>
+void BoolExpr() {
+ // expected-error at +1{{value of type 'NoBoolConversion' is not contextually convertible to 'bool'}}
+#pragma acc parallel self (NoBoolConversion{})
+ while(0);
+ // expected-error at +2{{no member named 'NotValid' in 'NoBoolConversion'}}
+ // expected-note@#INST{{in instantiation of function template specialization}}
+#pragma acc parallel self (T::NotValid)
+ while(0);
+
+#pragma acc parallel self (BoolConversion{})
+ while(0);
+
+ // expected-error at +1{{value of type 'NoBoolConversion' is not contextually convertible to 'bool'}}
+#pragma acc parallel self (T{})
+ while(0);
+
+#pragma acc parallel self (U{})
+ while(0);
+}
+
+struct HasBool {
+ static constexpr bool B = true;
+};
+
+template<typename T>
+void WarnMaybeNotUsed() {
+ // expected-warning at +2{{OpenACC construct 'self' has no effect when an 'if' clause evaluates to true}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc parallel self if(T::B)
+ while(0);
+
+ // expected-warning at +2{{OpenACC construct 'self' has no effect when an 'if' clause evaluates to true}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc parallel self(T::B) if(T::B)
+ while(0);
+
+ // expected-warning at +2{{OpenACC construct 'self' has no effect when an 'if' clause evaluates to true}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc parallel if(T::B) self
+ while(0);
+
+ // expected-warning at +2{{OpenACC construct 'self' has no effect when an 'if' clause evaluates to true}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc parallel if(T::B) self(T::B)
+ while(0);
+
+ // We still warn in the cases of dependent failures, since the diagnostic
+ // happens immediately rather than during instantiation.
+
+ // expected-error at +4{{no member named 'Invalid' in 'HasBool'}}
+ // expected-note@#NOT_USED_INST{{in instantiation of function template specialization 'WarnMaybeNotUsed<HasBool>' requested here}}
+ // expected-warning at +2{{OpenACC construct 'self' has no effect when an 'if' clause evaluates to true}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc parallel self if(T::Invalid)
+ while(0);
+
+ // expected-error at +3{{no member named 'Invalid' in 'HasBool'}}
+ // expected-warning at +2{{OpenACC construct 'self' has no effect when an 'if' clause evaluates to true}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc parallel self(T::Invalid) if(T::B)
+ while(0);
+
+ // expected-error at +3{{no member named 'Invalid' in 'HasBool'}}
+ // expected-warning at +2{{OpenACC construct 'self' has no effect when an 'if' clause evaluates to true}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc parallel self(T::B) if(T::Invalid)
+ while(0);
+
+ // expected-error at +3{{no member named 'Invalid' in 'HasBool'}}
+ // expected-warning at +2{{OpenACC construct 'self' has no effect when an 'if' clause evaluates to true}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc parallel if(T::Invalid) self
+ while(0);
+
+ // expected-error at +3{{no member named 'Invalid' in 'HasBool'}}
+ // expected-warning at +2{{OpenACC construct 'self' has no effect when an 'if' clause evaluates to true}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc parallel if(T::Invalid) self(T::B)
+ while(0);
+
+ // expected-error at +3{{no member named 'Invalid' in 'HasBool'}}
+ // expected-warning at +2{{OpenACC construct 'self' has no effect when an 'if' clause evaluates to true}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc parallel if(T::B) self(T::Invalid)
+ while(0);
+}
+
+void Instantiate() {
+ BoolExpr<NoBoolConversion, BoolConversion>(); // #INST
+ WarnMaybeNotUsed<HasBool>(); // #NOT_USED_INST
+}
diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
index f304786ff9dffd..2ef599d2cd26fa 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -2791,6 +2791,10 @@ void OpenACCClauseEnqueue::VisitDefaultClause(const OpenACCDefaultClause &C) {}
void OpenACCClauseEnqueue::VisitIfClause(const OpenACCIfClause &C) {
Visitor.AddStmt(C.getConditionExpr());
}
+void OpenACCClauseEnqueue::VisitSelfClause(const OpenACCSelfClause &C) {
+ if (C.hasConditionExpr())
+ Visitor.AddStmt(C.getConditionExpr());
+}
} // namespace
void EnqueueVisitor::EnqueueChildren(const OpenACCClause *C) {
More information about the cfe-commits
mailing list