https://github.com/erichkeane created
https://github.com/llvm/llvm-project/pull/112206
The worker clause specifies iterations of the loop/ that are executed in
parallel by distributing the iterations among the multiple works within a
single gang.
The sema rules for this type are simply that it cannot be combined with a
`kernel` construct with a `num_workers` clause, child `loop` clauses cannot
contain a `gang` or `worker` clause, and that the argument is oly allowed when
associated with a `kernel`.
>From bc9b06006b15a4f7f3bd4b89d9a5191a43280e04 Mon Sep 17 00:00:00 2001
From: erichkeane <eke...@nvidia.com>
Date: Fri, 11 Oct 2024 09:59:24 -0700
Subject: [PATCH] [OpenACC] implement loop 'worker' clause.
The worker clause specifies iterations of the loop/ that are executed in
parallel by distributing the iterations among the multiple works within
a single gang.
The sema rules for this type are simply that it cannot be combined with
a `kerne` construct with a `num_workers` clause, child `loop` clauses
cannot contain a `gang` or `worker` clause, and that the argument is oly
allowed when associated with a `kernel`.
---
clang/include/clang/AST/OpenACCClause.h | 42 ++-
.../clang/Basic/DiagnosticSemaKinds.td | 21 +-
clang/include/clang/Basic/OpenACCClauses.def | 1 +
clang/include/clang/Sema/SemaOpenACC.h | 16 ++
clang/lib/AST/OpenACCClause.cpp | 30 +-
clang/lib/AST/StmtProfile.cpp | 6 +
clang/lib/AST/TextNodeDumper.cpp | 1 +
clang/lib/Parse/ParseOpenACC.cpp | 1 +
clang/lib/Sema/SemaOpenACC.cpp | 125 +++++++-
clang/lib/Sema/TreeTransform.h | 28 ++
clang/lib/Serialization/ASTReader.cpp | 7 +-
clang/lib/Serialization/ASTWriter.cpp | 9 +-
.../AST/ast-print-openacc-loop-construct.cpp | 39 +++
clang/test/ParserOpenACC/parse-clauses.c | 19 +-
.../compute-construct-device_type-clause.c | 3 +-
...p-construct-auto_seq_independent-clauses.c | 15 +-
.../loop-construct-device_type-clause.c | 1 -
.../SemaOpenACC/loop-construct-worker-ast.cpp | 270 ++++++++++++++++++
.../loop-construct-worker-clause.cpp | 202 +++++++++++++
clang/tools/libclang/CIndex.cpp | 5 +
20 files changed, 766 insertions(+), 75 deletions(-)
create mode 100644 clang/test/SemaOpenACC/loop-construct-worker-ast.cpp
create mode 100644 clang/test/SemaOpenACC/loop-construct-worker-clause.cpp
diff --git a/clang/include/clang/AST/OpenACCClause.h
b/clang/include/clang/AST/OpenACCClause.h
index f3a09eb651458d..e8b8f477f91ae7 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -145,32 +145,6 @@ class OpenACCVectorClause : 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 OpenACCWorkerClause : public OpenACCClause {
-protected:
- OpenACCWorkerClause(SourceLocation BeginLoc, SourceLocation EndLoc)
- : OpenACCClause(OpenACCClauseKind::Worker, BeginLoc, EndLoc) {
- llvm_unreachable("Not yet implemented");
- }
-
-public:
- static bool classof(const OpenACCClause *C) {
- return C->getClauseKind() == OpenACCClauseKind::Worker;
- }
-
- static OpenACCWorkerClause *
- 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());
- }
-};
-
/// Represents a clause that has a list of parameters.
class OpenACCClauseWithParams : public OpenACCClause {
/// Location of the '('.
@@ -541,6 +515,22 @@ class OpenACCGangClause final
ArrayRef<Expr *> IntExprs, SourceLocation EndLoc);
};
+class OpenACCWorkerClause : public OpenACCClauseWithSingleIntExpr {
+protected:
+ OpenACCWorkerClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
+ Expr *IntExpr, SourceLocation EndLoc);
+
+public:
+ static bool classof(const OpenACCClause *C) {
+ return C->getClauseKind() == OpenACCClauseKind::Worker;
+ }
+
+ static OpenACCWorkerClause *Create(const ASTContext &Ctx,
+ SourceLocation BeginLoc,
+ SourceLocation LParenLoc, Expr *IntExpr,
+ 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 3c62a017005e59..00201a7f192f17 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12686,21 +12686,22 @@ def err_acc_intervening_code
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 "
+def err_acc_int_arg_invalid
+ : Error<"'%1' argument on '%0' clause is not permitted on a%select{n "
+ "orphaned|||}2 'loop' construct %select{|associated with a "
"'parallel' compute construct|associated with a 'kernels' compute "
- "construct|associated with a 'serial' compute construct}1">;
+ "construct|associated with a 'serial' compute construct}2">;
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' "
+def err_acc_num_arg_conflict
+ : Error<"'num' argument to '%0' clause not allowed on a 'loop' construct "
+ "associated with a 'kernels' construct that has a "
+ "'%select{num_gangs|num_workers}1' "
"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">;
+def err_acc_clause_in_clause_region
+ : Error<"loop with a '%0' clause may not exist in the region of a '%1' "
+ "clause%select{| on a 'kernels' compute construct}2">;
// 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 2a098de31eb618..4c0b56dc13e625 100644
--- a/clang/include/clang/Basic/OpenACCClauses.def
+++ b/clang/include/clang/Basic/OpenACCClauses.def
@@ -56,6 +56,7 @@ VISIT_CLAUSE(Seq)
VISIT_CLAUSE(Tile)
VISIT_CLAUSE(VectorLength)
VISIT_CLAUSE(Wait)
+VISIT_CLAUSE(Worker)
#undef VISIT_CLAUSE
#undef CLAUSE_ALIAS
diff --git a/clang/include/clang/Sema/SemaOpenACC.h
b/clang/include/clang/Sema/SemaOpenACC.h
index 59a9648d5f9380..e253610a84b0bf 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -118,6 +118,11 @@ class SemaOpenACC : public SemaBase {
/// 'kernel' construct, this will have the source location for it. This
/// permits us to implement the restriction of no further 'gang' clauses.
SourceLocation LoopGangClauseOnKernelLoc;
+ /// If there is a current 'active' loop construct with a 'worker' clause on
it
+ /// (on any sort of construct), this has the source location for it. This
+ /// permits us to implement the restriction of no further 'gang' or 'worker'
+ /// clauses.
+ SourceLocation LoopWorkerClauseLoc;
// Redeclaration of the version in OpenACCClause.h.
using DeviceTypeArgument = std::pair<IdentifierInfo *, SourceLocation>;
@@ -224,11 +229,15 @@ class SemaOpenACC : public SemaBase {
ClauseKind == OpenACCClauseKind::NumWorkers ||
ClauseKind == OpenACCClauseKind::Async ||
ClauseKind == OpenACCClauseKind::Tile ||
+ ClauseKind == OpenACCClauseKind::Worker ||
+ ClauseKind == OpenACCClauseKind::Vector ||
ClauseKind == OpenACCClauseKind::VectorLength) &&
"Parsed clause kind does not have a int exprs");
// 'async' and 'wait' have an optional IntExpr, so be tolerant of that.
if ((ClauseKind == OpenACCClauseKind::Async ||
+ ClauseKind == OpenACCClauseKind::Worker ||
+ ClauseKind == OpenACCClauseKind::Vector ||
ClauseKind == OpenACCClauseKind::Wait) &&
std::holds_alternative<std::monostate>(Details))
return 0;
@@ -271,6 +280,8 @@ class SemaOpenACC : public SemaBase {
ClauseKind == OpenACCClauseKind::Async ||
ClauseKind == OpenACCClauseKind::Tile ||
ClauseKind == OpenACCClauseKind::Gang ||
+ ClauseKind == OpenACCClauseKind::Worker ||
+ ClauseKind == OpenACCClauseKind::Vector ||
ClauseKind == OpenACCClauseKind::VectorLength) &&
"Parsed clause kind does not have a int exprs");
@@ -401,6 +412,8 @@ class SemaOpenACC : public SemaBase {
ClauseKind == OpenACCClauseKind::NumWorkers ||
ClauseKind == OpenACCClauseKind::Async ||
ClauseKind == OpenACCClauseKind::Tile ||
+ ClauseKind == OpenACCClauseKind::Worker ||
+ ClauseKind == OpenACCClauseKind::Vector ||
ClauseKind == OpenACCClauseKind::VectorLength) &&
"Parsed clause kind does not have a int exprs");
Details = IntExprDetails{{IntExprs.begin(), IntExprs.end()}};
@@ -410,6 +423,8 @@ class SemaOpenACC : public SemaBase {
ClauseKind == OpenACCClauseKind::NumWorkers ||
ClauseKind == OpenACCClauseKind::Async ||
ClauseKind == OpenACCClauseKind::Tile ||
+ ClauseKind == OpenACCClauseKind::Worker ||
+ ClauseKind == OpenACCClauseKind::Vector ||
ClauseKind == OpenACCClauseKind::VectorLength) &&
"Parsed clause kind does not have a int exprs");
Details = IntExprDetails{std::move(IntExprs)};
@@ -663,6 +678,7 @@ class SemaOpenACC : public SemaBase {
ComputeConstructInfo OldActiveComputeConstructInfo;
OpenACCDirectiveKind DirKind;
SourceLocation OldLoopGangClauseOnKernelLoc;
+ SourceLocation OldLoopWorkerClauseLoc;
llvm::SmallVector<OpenACCLoopConstruct *> ParentlessLoopConstructs;
LoopInConstructRAII LoopRAII;
diff --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp
index 6fb8fe0b8cfeef..638252fd811f1d 100644
--- a/clang/lib/AST/OpenACCClause.cpp
+++ b/clang/lib/AST/OpenACCClause.cpp
@@ -44,7 +44,8 @@ bool OpenACCClauseWithCondition::classof(const OpenACCClause
*C) {
bool OpenACCClauseWithSingleIntExpr::classof(const OpenACCClause *C) {
return OpenACCNumWorkersClause::classof(C) ||
OpenACCVectorLengthClause::classof(C) ||
- OpenACCCollapseClause::classof(C) || OpenACCAsyncClause::classof(C);
+ OpenACCWorkerClause::classof(C) || OpenACCCollapseClause::classof(C)
||
+ OpenACCAsyncClause::classof(C);
}
OpenACCDefaultClause *OpenACCDefaultClause::Create(const ASTContext &C,
OpenACCDefaultClauseKind K,
@@ -403,11 +404,24 @@ OpenACCGangClause::Create(const ASTContext &C,
SourceLocation BeginLoc,
OpenACCGangClause(BeginLoc, LParenLoc, GangKinds, IntExprs, EndLoc);
}
+OpenACCWorkerClause::OpenACCWorkerClause(SourceLocation BeginLoc,
+ SourceLocation LParenLoc,
+ Expr *IntExpr, SourceLocation EndLoc)
+ : OpenACCClauseWithSingleIntExpr(OpenACCClauseKind::Worker, BeginLoc,
+ LParenLoc, IntExpr, EndLoc) {
+ assert((!IntExpr || IntExpr->isInstantiationDependent() ||
+ IntExpr->getType()->isIntegerType()) &&
+ "Int expression type not scalar/dependent");
+}
+
OpenACCWorkerClause *OpenACCWorkerClause::Create(const ASTContext &C,
SourceLocation BeginLoc,
+ SourceLocation LParenLoc,
+ Expr *IntExpr,
SourceLocation EndLoc) {
- void *Mem = C.Allocate(sizeof(OpenACCWorkerClause));
- return new (Mem) OpenACCWorkerClause(BeginLoc, EndLoc);
+ void *Mem =
+ C.Allocate(sizeof(OpenACCWorkerClause), alignof(OpenACCWorkerClause));
+ return new (Mem) OpenACCWorkerClause(BeginLoc, LParenLoc, IntExpr, EndLoc);
}
OpenACCVectorClause *OpenACCVectorClause::Create(const ASTContext &C,
@@ -638,3 +652,13 @@ void OpenACCClausePrinter::VisitGangClause(const
OpenACCGangClause &C) {
OS << ")";
}
}
+
+void OpenACCClausePrinter::VisitWorkerClause(const OpenACCWorkerClause &C) {
+ OS << "worker";
+
+ if (C.hasIntExpr()) {
+ OS << "(num: ";
+ printExpr(C.getIntExpr());
+ OS << ")";
+ }
+}
diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index 6161b1403ed35d..25b1cbb8590869 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -2629,6 +2629,12 @@ void OpenACCClauseProfiler::VisitAsyncClause(const
OpenACCAsyncClause &Clause) {
Profiler.VisitStmt(Clause.getIntExpr());
}
+void OpenACCClauseProfiler::VisitWorkerClause(
+ const OpenACCWorkerClause &Clause) {
+ if (Clause.hasIntExpr())
+ Profiler.VisitStmt(Clause.getIntExpr());
+}
+
void OpenACCClauseProfiler::VisitWaitClause(const OpenACCWaitClause &Clause) {
if (Clause.hasDevNumExpr())
Profiler.VisitStmt(Clause.getDevNumExpr());
diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp
index ac8c196777f9b8..beccb0615f0e9c 100644
--- a/clang/lib/AST/TextNodeDumper.cpp
+++ b/clang/lib/AST/TextNodeDumper.cpp
@@ -420,6 +420,7 @@ void TextNodeDumper::Visit(const OpenACCClause *C) {
case OpenACCClauseKind::Self:
case OpenACCClauseKind::Seq:
case OpenACCClauseKind::Tile:
+ case OpenACCClauseKind::Worker:
case OpenACCClauseKind::VectorLength:
// 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.
diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp
index 635039b724e6a0..51d4dc38c17f67 100644
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@@ -1130,6 +1130,7 @@ Parser::OpenACCClauseParseResult
Parser::ParseOpenACCClauseParams(
Parens.skipToEnd();
return OpenACCCanContinue();
}
+ ParsedClause.setIntExprDetails(IntExpr.get());
break;
}
case OpenACCClauseKind::Async: {
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 30d73d621db69b..1b24331cbd87ca 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -377,6 +377,18 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind
DirectiveKind,
default:
return false;
}
+ case OpenACCClauseKind::Worker: {
+ switch (DirectiveKind) {
+ case OpenACCDirectiveKind::Loop:
+ case OpenACCDirectiveKind::ParallelLoop:
+ case OpenACCDirectiveKind::SerialLoop:
+ case OpenACCDirectiveKind::KernelsLoop:
+ case OpenACCDirectiveKind::Routine:
+ return true;
+ default:
+ return false;
+ }
+ }
}
default:
@@ -500,7 +512,6 @@ class SemaOpenACCClauseVisitor {
OpenACCClause *Visit(SemaOpenACC::OpenACCParsedClause &Clause) {
switch (Clause.getClauseKind()) {
- case OpenACCClauseKind::Worker:
case OpenACCClauseKind::Vector: {
// TODO OpenACC: These are only implemented enough for the 'seq'
// diagnostic, otherwise treats itself as unimplemented. When we
@@ -1024,6 +1035,75 @@ OpenACCClause
*SemaOpenACCClauseVisitor::VisitIndependentClause(
Clause.getEndLoc());
}
+OpenACCClause *SemaOpenACCClauseVisitor::VisitWorkerClause(
+ 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();
+
+ Expr *IntExpr =
+ Clause.getNumIntExprs() != 0 ? Clause.getIntExprs()[0] : nullptr;
+
+ if (IntExpr) {
+ switch (SemaRef.getActiveComputeConstructInfo().Kind) {
+ case OpenACCDirectiveKind::Invalid:
+ SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_int_arg_invalid)
+ << OpenACCClauseKind::Worker << "num" << /*orphan=*/0;
+ IntExpr = nullptr;
+ break;
+ case OpenACCDirectiveKind::Parallel:
+ SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_int_arg_invalid)
+ << OpenACCClauseKind::Worker << "num" << /*parallel=*/1;
+ IntExpr = nullptr;
+ break;
+ case OpenACCDirectiveKind::Serial:
+ SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_int_arg_invalid)
+ << OpenACCClauseKind::Worker << "num" << /*serial=*/3;
+ IntExpr = nullptr;
+ break;
+ case OpenACCDirectiveKind::Kernels: {
+ const auto *Itr =
+ llvm::find_if(SemaRef.getActiveComputeConstructInfo().Clauses,
+ llvm::IsaPred<OpenACCNumWorkersClause>);
+ if (Itr != SemaRef.getActiveComputeConstructInfo().Clauses.end()) {
+ SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict)
+ << OpenACCClauseKind::Worker << /*num_workers=*/1;
+ SemaRef.Diag((*Itr)->getBeginLoc(),
+ diag::note_acc_previous_clause_here);
+
+ IntExpr = nullptr;
+ }
+ break;
+ }
+ default:
+ llvm_unreachable("Non compute construct in active compute construct");
+ }
+ }
+
+ // OpenACC 3.3 2.9.3: The region of a loop with a 'worker' clause may not
+ // contain a loop with a gang or worker clause unless within a nested compute
+ // region.
+ if (SemaRef.LoopWorkerClauseLoc.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_clause_in_clause_region)
+ << OpenACCClauseKind::Worker << OpenACCClauseKind::Worker
+ << /*skip kernels construct info*/ 0;
+ SemaRef.Diag(SemaRef.LoopWorkerClauseLoc,
+ diag::note_acc_previous_clause_here);
+ return nullptr;
+ }
+
+ return OpenACCWorkerClause::Create(Ctx, Clause.getBeginLoc(),
+ Clause.getLParenLoc(), IntExpr,
+ Clause.getEndLoc());
+}
+
OpenACCClause *SemaOpenACCClauseVisitor::VisitGangClause(
SemaOpenACC::OpenACCParsedClause &Clause) {
if (DiagIfSeqClause(Clause))
@@ -1061,8 +1141,8 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitGangClause(
llvm::IsaPred<OpenACCNumGangsClause>);
if (Itr != SemaRef.getActiveComputeConstructInfo().Clauses.end()) {
- SemaRef.Diag(ER.get()->getBeginLoc(),
- diag::err_acc_gang_num_gangs_conflict);
+ SemaRef.Diag(ER.get()->getBeginLoc(), diag::err_acc_num_arg_conflict)
+ << OpenACCClauseKind::Gang << /*num_gangs=*/0;
SemaRef.Diag((*Itr)->getBeginLoc(),
diag::note_acc_previous_clause_here);
continue;
@@ -1091,12 +1171,28 @@ OpenACCClause
*SemaOpenACCClauseVisitor::VisitGangClause(
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(Clause.getBeginLoc(), diag::err_acc_clause_in_clause_region)
+ << OpenACCClauseKind::Gang << OpenACCClauseKind::Gang
+ << /*kernels construct info*/ 1;
SemaRef.Diag(SemaRef.LoopGangClauseOnKernelLoc,
diag::note_acc_previous_clause_here);
return nullptr;
}
+ // OpenACC 3.3 2.9.3: The region of a loop with a 'worker' clause may not
+ // contain a loop with a gang or worker clause unless within a nested compute
+ // region.
+ if (SemaRef.LoopWorkerClauseLoc.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_clause_in_clause_region)
+ << OpenACCClauseKind::Gang << OpenACCClauseKind::Worker
+ << /*kernels construct info*/ 1;
+ SemaRef.Diag(SemaRef.LoopWorkerClauseLoc,
+ diag::note_acc_previous_clause_here);
+ return nullptr;
+ }
+
return OpenACCGangClause::Create(Ctx, Clause.getBeginLoc(),
Clause.getLParenLoc(), GangKinds, IntExprs,
Clause.getEndLoc());
@@ -1216,6 +1312,7 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
ArrayRef<OpenACCClause *> Clauses)
: SemaRef(S), OldActiveComputeConstructInfo(S.ActiveComputeConstructInfo),
DirKind(DK), OldLoopGangClauseOnKernelLoc(S.LoopGangClauseOnKernelLoc),
+ OldLoopWorkerClauseLoc(S.LoopWorkerClauseLoc),
LoopRAII(SemaRef, /*PreserveDepth=*/false) {
// Compute constructs end up taking their 'loop'.
if (DirKind == OpenACCDirectiveKind::Parallel ||
@@ -1232,6 +1329,7 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
//
// Implement the 'unless within a nested compute region' part.
SemaRef.LoopGangClauseOnKernelLoc = {};
+ SemaRef.LoopWorkerClauseLoc = {};
} else if (DirKind == OpenACCDirectiveKind::Loop) {
SetCollapseInfoBeforeAssociatedStmt(UnInstClauses, Clauses);
SetTileInfoBeforeAssociatedStmt(UnInstClauses, Clauses);
@@ -1252,6 +1350,12 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
if (Itr != Clauses.end())
SemaRef.LoopGangClauseOnKernelLoc = (*Itr)->getBeginLoc();
}
+
+ if (UnInstClauses.empty()) {
+ auto *Itr = llvm::find_if(Clauses, llvm::IsaPred<OpenACCWorkerClause>);
+ if (Itr != Clauses.end())
+ SemaRef.LoopWorkerClauseLoc = (*Itr)->getBeginLoc();
+ }
}
}
@@ -1324,6 +1428,7 @@ void
SemaOpenACC::AssociatedStmtRAII::SetTileInfoBeforeAssociatedStmt(
SemaOpenACC::AssociatedStmtRAII::~AssociatedStmtRAII() {
SemaRef.ActiveComputeConstructInfo = OldActiveComputeConstructInfo;
SemaRef.LoopGangClauseOnKernelLoc = OldLoopGangClauseOnKernelLoc;
+ SemaRef.LoopWorkerClauseLoc = OldLoopWorkerClauseLoc;
if (DirKind == OpenACCDirectiveKind::Parallel ||
DirKind == OpenACCDirectiveKind::Serial ||
@@ -1934,8 +2039,8 @@ ExprResult SemaOpenACC::CheckGangExpr(OpenACCGangKind GK,
Expr *E) {
// 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
+ Diag(E->getBeginLoc(), diag::err_acc_int_arg_invalid)
+ << OpenACCClauseKind::Gang << GK
<< (/*orphan/parallel=*/ActiveComputeConstructInfo.Kind ==
OpenACCDirectiveKind::Parallel
? 1
@@ -1951,8 +2056,8 @@ ExprResult SemaOpenACC::CheckGangExpr(OpenACCGangKind GK,
Expr *E) {
// 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;
+ Diag(E->getBeginLoc(), diag::err_acc_int_arg_invalid)
+ << OpenACCClauseKind::Gang << 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
@@ -1975,8 +2080,8 @@ ExprResult SemaOpenACC::CheckGangExpr(OpenACCGangKind GK,
Expr *E) {
// too, so we disallow them too.
case OpenACCGangKind::Dim:
case OpenACCGangKind::Num:
- Diag(E->getBeginLoc(), diag::err_acc_gang_arg_invalid)
- << GK << /*Kernels=*/3;
+ Diag(E->getBeginLoc(), diag::err_acc_int_arg_invalid)
+ << OpenACCClauseKind::Gang << GK << /*Kernels=*/3;
return ExprError();
case OpenACCGangKind::Static:
return CheckGangStaticExpr(*this, E);
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index cde40773336866..45e8b3cf6bd8fc 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -11787,6 +11787,34 @@ void OpenACCClauseTransform<Derived>::VisitAsyncClause(
: nullptr,
ParsedClause.getEndLoc());
}
+
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitWorkerClause(
+ const OpenACCWorkerClause &C) {
+ if (C.hasIntExpr()) {
+ // restrictions on this expression are all "does it exist in certain
+ // situations" that are not possible to be dependent, so the only check we
+ // have is that it transforms, and is an int expression.
+ ExprResult Res = Self.TransformExpr(const_cast<Expr *>(C.getIntExpr()));
+ if (!Res.isUsable())
+ return;
+
+ Res = Self.getSema().OpenACC().ActOnIntExpr(OpenACCDirectiveKind::Invalid,
+ C.getClauseKind(),
+ C.getBeginLoc(), Res.get());
+ if (!Res.isUsable())
+ return;
+ ParsedClause.setIntExprDetails(Res.get());
+ }
+
+ NewClause = OpenACCWorkerClause::Create(
+ Self.getSema().getASTContext(), ParsedClause.getBeginLoc(),
+ ParsedClause.getLParenLoc(),
+ ParsedClause.getNumIntExprs() != 0 ? ParsedClause.getIntExprs()[0]
+ : nullptr,
+ ParsedClause.getEndLoc());
+}
+
template <typename Derived>
void OpenACCClauseTransform<Derived>::VisitWaitClause(
const OpenACCWaitClause &C) {
diff --git a/clang/lib/Serialization/ASTReader.cpp
b/clang/lib/Serialization/ASTReader.cpp
index 31881a39de47f7..ecc5d3c59a3549 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -12339,10 +12339,15 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
return OpenACCGangClause::Create(getContext(), BeginLoc, LParenLoc,
GangKinds, Exprs, EndLoc);
}
+ case OpenACCClauseKind::Worker: {
+ SourceLocation LParenLoc = readSourceLocation();
+ Expr *WorkerExpr = readBool() ? readSubExpr() : nullptr;
+ return OpenACCWorkerClause::Create(getContext(), BeginLoc, LParenLoc,
+ WorkerExpr, EndLoc);
+ }
case OpenACCClauseKind::Finalize:
case OpenACCClauseKind::IfPresent:
- case OpenACCClauseKind::Worker:
case OpenACCClauseKind::Vector:
case OpenACCClauseKind::NoHost:
case OpenACCClauseKind::UseDevice:
diff --git a/clang/lib/Serialization/ASTWriter.cpp
b/clang/lib/Serialization/ASTWriter.cpp
index 583d9a4bccb800..0a6e260e3e4e93 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -8192,10 +8192,17 @@ void ASTRecordWriter::writeOpenACCClause(const
OpenACCClause *C) {
}
return;
}
+ case OpenACCClauseKind::Worker: {
+ const auto *WC = cast<OpenACCWorkerClause>(C);
+ writeSourceLocation(WC->getLParenLoc());
+ writeBool(WC->hasIntExpr());
+ if (WC->hasIntExpr())
+ AddStmt(const_cast<Expr *>(WC->getIntExpr()));
+ return;
+ }
case OpenACCClauseKind::Finalize:
case OpenACCClauseKind::IfPresent:
- case OpenACCClauseKind::Worker:
case OpenACCClauseKind::Vector:
case OpenACCClauseKind::NoHost:
case OpenACCClauseKind::UseDevice:
diff --git a/clang/test/AST/ast-print-openacc-loop-construct.cpp
b/clang/test/AST/ast-print-openacc-loop-construct.cpp
index baa4b173f88edc..ee11435aaa4b1c 100644
--- a/clang/test/AST/ast-print-openacc-loop-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-loop-construct.cpp
@@ -177,4 +177,43 @@ void foo() {
#pragma acc serial
#pragma acc loop gang
for(;;);
+
+// CHECK: #pragma acc loop worker
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc parallel
+#pragma acc loop worker
+ for(;;);
+
+// CHECK: #pragma acc parallel
+// CHECK-NEXT: #pragma acc loop worker
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc parallel
+#pragma acc loop worker
+ for(;;);
+
+// CHECK: #pragma acc serial
+// CHECK-NEXT: #pragma acc loop worker
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc serial
+#pragma acc loop worker
+ for(;;);
+
+// CHECK: #pragma acc kernels
+// CHECK-NEXT: #pragma acc loop worker(num: 5)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc kernels
+#pragma acc loop worker(5)
+ for(;;);
+
+// CHECK: #pragma acc kernels
+// CHECK-NEXT: #pragma acc loop worker(num: 5)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc kernels
+#pragma acc loop worker(num:5)
+ for(;;);
}
diff --git a/clang/test/ParserOpenACC/parse-clauses.c
b/clang/test/ParserOpenACC/parse-clauses.c
index 899fbd78b87298..81c48335cf0c42 100644
--- a/clang/test/ParserOpenACC/parse-clauses.c
+++ b/clang/test/ParserOpenACC/parse-clauses.c
@@ -949,7 +949,6 @@ void IntExprParsing() {
#pragma acc loop vector(length:returns_int())
for(;;);
- // expected-warning@+1{{OpenACC clause 'worker' not yet implemented, clause
ignored}}
#pragma acc loop worker
for(;;);
// expected-error@+1{{expected expression}}
@@ -959,8 +958,8 @@ void IntExprParsing() {
// expected-error@+1{{expected expression}}
#pragma acc loop worker(invalid:)
for(;;);
- // expected-error@+2{{invalid tag 'invalid' on 'worker' clause}}
- // expected-warning@+1{{OpenACC clause 'worker' not yet implemented, clause
ignored}}
+#pragma acc kernels
+ // expected-error@+1{{invalid tag 'invalid' on 'worker' clause}}
#pragma acc loop worker(invalid:5)
for(;;);
// expected-error@+1{{expected expression}}
@@ -983,21 +982,21 @@ void IntExprParsing() {
// expected-note@+1{{to match this '('}}
#pragma acc loop worker(length:6,4)
for(;;);
- // expected-warning@+1{{OpenACC clause 'worker' not yet implemented, clause
ignored}}
+#pragma acc kernels
#pragma acc loop worker(5)
for(;;);
- // expected-error@+2{{invalid tag 'length' on 'worker' clause}}
- // expected-warning@+1{{OpenACC clause 'worker' not yet implemented, clause
ignored}}
+#pragma acc kernels
+ // expected-error@+1{{invalid tag 'length' on 'worker' clause}}
#pragma acc loop worker(length:5)
for(;;);
- // expected-warning@+1{{OpenACC clause 'worker' not yet implemented, clause
ignored}}
+#pragma acc kernels
#pragma acc loop worker(num:5)
for(;;);
- // expected-warning@+1{{OpenACC clause 'worker' not yet implemented, clause
ignored}}
+#pragma acc kernels
#pragma acc loop worker(returns_int())
for(;;);
- // expected-error@+2{{invalid tag 'length' on 'worker' clause}}
- // expected-warning@+1{{OpenACC clause 'worker' not yet implemented, clause
ignored}}
+#pragma acc kernels
+ // expected-error@+1{{invalid tag 'length' on 'worker' clause}}
#pragma acc loop worker(length:returns_int())
for(;;);
}
diff --git a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
index 89000517c43fb5..aaf8b76e1f3dfb 100644
--- a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
@@ -61,8 +61,7 @@ void uses() {
// expected-error@+1{{OpenACC 'auto' clause is not valid on 'kernels'
directive}}
#pragma acc kernels device_type(*) auto
while(1);
- // expected-error@+2{{OpenACC clause 'worker' may not follow a 'device_type'
clause in a compute construct}}
- // expected-note@+1{{previous clause is here}}
+ // expected-error@+1{{OpenACC 'worker' clause is not valid on 'kernels'
directive}}
#pragma acc kernels device_type(*) worker
while(1);
// expected-error@+2{{OpenACC clause 'nohost' may not follow a 'device_type'
clause in a compute construct}}
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 6c2c79b02a4131..6a975956f3ff5c 100644
--- a/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
+++ b/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
@@ -43,7 +43,6 @@ void uses() {
// expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}}
#pragma acc loop auto if_present
for(;;);
- // expected-warning@+1{{OpenACC clause 'worker' not yet implemented}}
#pragma acc loop auto worker
for(;;);
// expected-warning@+1{{OpenACC clause 'vector' not yet implemented}}
@@ -180,7 +179,6 @@ void uses() {
// expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}}
#pragma acc loop if_present auto
for(;;);
- // expected-warning@+1{{OpenACC clause 'worker' not yet implemented}}
#pragma acc loop worker auto
for(;;);
// expected-warning@+1{{OpenACC clause 'vector' not yet implemented}}
@@ -318,7 +316,6 @@ void uses() {
// expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}}
#pragma acc loop independent if_present
for(;;);
- // expected-warning@+1{{OpenACC clause 'worker' not yet implemented}}
#pragma acc loop independent worker
for(;;);
// expected-warning@+1{{OpenACC clause 'vector' not yet implemented}}
@@ -455,7 +452,6 @@ void uses() {
// expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}}
#pragma acc loop if_present independent
for(;;);
- // expected-warning@+1{{OpenACC clause 'worker' not yet implemented}}
#pragma acc loop worker independent
for(;;);
// expected-warning@+1{{OpenACC clause 'vector' not yet implemented}}
@@ -591,9 +587,8 @@ void uses() {
// expected-note@+1{{previous clause is here}}
#pragma acc loop seq gang
for(;;);
- // expected-error@+3{{OpenACC clause 'worker' may not appear on the same
construct as a 'seq' clause on a 'loop' construct}}
- // expected-note@+2{{previous clause is here}}
- // expected-warning@+1{{OpenACC clause 'worker' not yet implemented}}
+ // expected-error@+2{{OpenACC clause 'worker' may not appear on the same
construct as a 'seq' clause on a 'loop' construct}}
+ // expected-note@+1{{previous clause is here}}
#pragma acc loop seq worker
for(;;);
// expected-error@+3{{OpenACC clause 'vector' may not appear on the same
construct as a 'seq' clause on a 'loop' construct}}
@@ -734,10 +729,8 @@ void uses() {
// expected-note@+1{{previous clause is here}}
#pragma acc loop gang seq
for(;;);
- // TODO OpenACC: when 'worker' is implemented and makes it to the AST, this
should diagnose because of a conflict with 'seq'.
- // TODOexpected-error@+3{{OpenACC clause 'worker' may not appear on the same
construct as a 'seq' clause on a 'loop' construct}}
- // TODOexpected-note@+2{{previous clause is here}}
- // expected-warning@+1{{OpenACC clause 'worker' not yet implemented}}
+ // expected-error@+2{{OpenACC clause 'seq' may not appear on the same
construct as a 'worker' clause on a 'loop' construct}}
+ // expected-note@+1{{previous clause is here}}
#pragma acc loop worker seq
for(;;);
// TODO OpenACC: when 'vector' 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 cedef3ca858f5e..51da8565f4e399 100644
--- a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
@@ -56,7 +56,6 @@ void uses() {
for(;;);
#pragma acc loop device_type(*) auto
for(;;);
- // expected-warning@+1{{OpenACC clause 'worker' not yet implemented, clause
ignored}}
#pragma acc loop device_type(*) worker
for(;;);
// expected-error@+2{{OpenACC clause 'nohost' may not follow a 'device_type'
clause in a 'loop' construct}}
diff --git a/clang/test/SemaOpenACC/loop-construct-worker-ast.cpp
b/clang/test/SemaOpenACC/loop-construct-worker-ast.cpp
new file mode 100644
index 00000000000000..6347e1419fd5c6
--- /dev/null
+++ b/clang/test/SemaOpenACC/loop-construct-worker-ast.cpp
@@ -0,0 +1,270 @@
+// 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
+
+template<unsigned I, typename ConvertsToInt, typename Int>
+void TemplUses(ConvertsToInt CTI, Int IsI) {
+ // CHECK: FunctionTemplateDecl{{.*}}TemplUses
+ // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'unsigned int' depth 0 index 0 I
+ // CHECK-NEXT: TemplateTypeParmDecl{{.*}}typename depth 0 index 1
ConvertsToInt
+ // CHECK-NEXT: TemplateTypeParmDecl{{.*}}typename depth 0 index 2 Int
+ // CHECK-NEXT: FunctionDecl{{.*}}TemplUses 'void (ConvertsToInt, Int)'
+ // CHECK-NEXT: ParmVarDecl{{.*}}CTI 'ConvertsToInt'
+ // CHECK-NEXT: ParmVarDecl{{.*}}IsI 'Int'
+ // CHECK-NEXT: CompoundStmt
+
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: worker clause{{.*}}
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+#pragma acc loop worker
+ for(;;);
+
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}}
parallel
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: worker clause{{.*}}
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+#pragma acc parallel
+#pragma acc loop worker
+ for(;;);
+
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}}
serial
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: worker 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 worker
+ for(;;);
+
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}}
kernels
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: worker clause{{.*}}
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'ConvertsToInt' lvalue ParmVar
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+#pragma acc kernels
+#pragma acc loop worker(CTI)
+ for(;;);
+
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}}
kernels
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: worker clause{{.*}}
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'Int' lvalue ParmVar
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+#pragma acc kernels
+#pragma acc loop worker(num:IsI)
+ for(;;);
+
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}}
kernels
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: worker clause{{.*}}
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'unsigned int' NonTypeTemplateParm{{.*}}'I'
'unsigned 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 worker(num:I)
+ for(;;);
+
+ // Instantiations:
+ // CHECK-NEXT: FunctionDecl{{.*}} TemplUses 'void (Converts, int)'
implicit_instantiation
+ // CHECK-NEXT: TemplateArgument integral '3U'
+ // CHECK-NEXT: TemplateArgument type 'Converts'
+ // CHECK-NEXT: RecordType{{.*}}'Converts'
+ // CHECK-NEXT: CXXRecord{{.*}}'Converts
+ // CHECK-NEXT: TemplateArgument type 'int'
+ // CHECK-NEXT: BuiltinType{{.*}}'int'
+ // CHECK-NEXT: ParmVarDecl{{.*}} CTI 'Converts'
+ // CHECK-NEXT: ParmVarDecl{{.*}} IsI 'int'
+ // CHECK-NEXT: CompoundStmt
+ //
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: worker clause{{.*}}
+ // 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: worker clause{{.*}}
+ // 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: worker clause{{.*}}
+ // 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]+]]{{.*}}
kernels
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: worker clause{{.*}}
+ // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <UserDefinedConversion>
+ // CHECK-NEXT: CXXMemberCallExpr{{.*}} 'int'
+ // CHECK-NEXT: MemberExpr{{.*}} .operator int
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'Converts' lvalue ParmVar
+ // 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]+]]{{.*}}
kernels
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: worker clause{{.*}}
+ // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar{{.*}} 'IsI' '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]+]]{{.*}}
kernels
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: worker clause{{.*}}
+ // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}}'unsigned int'
+ // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}}'unsigned int' depth 0 index 0 I
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'unsigned int' 3
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+}
+
+struct Converts{
+ operator int();
+};
+
+void uses() {
+ // CHECK: FunctionDecl{{.*}} uses
+ // CHECK-NEXT: CompoundStmt
+ //
+ // CHECK-NEXT: CallExpr
+ TemplUses<3>(Converts{}, 5);
+
+ // CHECK: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: worker clause{{.*}}
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+#pragma acc loop worker
+ for(;;);
+
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}}
parallel
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: worker clause{{.*}}
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+#pragma acc parallel
+#pragma acc loop worker
+ for(;;);
+
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}}
serial
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: worker 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 worker
+ for(;;);
+
+ Converts CTI;
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl
+ // CHECK-NEXT: CXXConstructExpr
+
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}}
kernels
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: worker clause{{.*}}
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <UserDefinedConversion>
+ // CHECK-NEXT: CXXMemberCallExpr{{.*}} 'int'
+ // CHECK-NEXT: MemberExpr{{.*}} .operator int
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Converts' lvalue Var
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+#pragma acc kernels
+#pragma acc loop worker(CTI)
+ for(;;);
+
+ int IsI;
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl
+
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}}
kernels
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: worker clause{{.*}}
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+ // CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue Var
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+#pragma acc kernels
+#pragma acc loop worker(num:IsI)
+ for(;;);
+}
+
+#endif // PCH_HELPER
diff --git a/clang/test/SemaOpenACC/loop-construct-worker-clause.cpp
b/clang/test/SemaOpenACC/loop-construct-worker-clause.cpp
new file mode 100644
index 00000000000000..f7d2f365a5aa79
--- /dev/null
+++ b/clang/test/SemaOpenACC/loop-construct-worker-clause.cpp
@@ -0,0 +1,202 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+template<unsigned I, typename NotInt, typename ConvertsToInt, typename Int>
+void TemplUses(NotInt NI, ConvertsToInt CTI, Int IsI) {
+ int i;
+
+ // expected-error@+1{{'num' argument on 'worker' clause is not permitted on
an orphaned 'loop' construct}}
+#pragma acc loop worker(i)
+ for(;;);
+
+ // expected-error@+1{{'num' argument on 'worker' clause is not permitted on
an orphaned 'loop' construct}}
+#pragma acc loop worker(num:IsI)
+ for(;;);
+
+#pragma acc kernels
+#pragma acc loop worker
+ for(;;);
+
+#pragma acc kernels
+#pragma acc loop worker(i)
+ for(;;);
+
+#pragma acc kernels
+#pragma acc loop worker(CTI)
+ for(;;);
+
+#pragma acc kernels
+#pragma acc loop worker(IsI)
+ for(;;);
+
+#pragma acc kernels
+#pragma acc loop worker(I)
+ for(;;);
+
+#pragma acc kernels
+ // expected-error@+1{{OpenACC clause 'worker' requires expression of integer
type ('NoConvert' invalid)}}
+#pragma acc loop worker(NI)
+ for(;;);
+
+#pragma acc kernels
+#pragma acc loop worker(num:i)
+ for(;;);
+
+ // expected-error@+3{{'num' argument to 'worker' clause not allowed on a
'loop' construct associated with a 'kernels' construct that has a 'num_workers'
clause}}
+ // expected-note@+1{{previous clause is here}}
+#pragma acc kernels num_workers(IsI)
+#pragma acc loop worker(num:CTI)
+ for(;;);
+ for(;;);
+}
+
+struct NoConvert{};
+struct Converts{
+ operator int();
+};
+
+void uses() {
+ TemplUses<3>(NoConvert{}, Converts{}, 5); // expected-note{{in instantiation
of function template specialization}}
+
+#pragma acc loop worker
+ for(;;);
+
+#pragma acc parallel
+#pragma acc loop worker
+ for(;;);
+
+ int i;
+
+ // expected-error@+1{{'num' argument on 'worker' clause is not permitted on
an orphaned 'loop' construct}}
+#pragma acc loop worker(i)
+ for(;;);
+
+ // expected-error@+2{{'num' argument on 'worker' clause is not permitted on
a 'loop' construct associated with a 'parallel' compute construct}}
+#pragma acc parallel
+#pragma acc loop worker(i)
+ for(;;);
+
+ // expected-error@+1{{'num' argument on 'worker' clause is not permitted on
an orphaned 'loop' construct}}
+#pragma acc loop worker(num:i)
+ for(;;);
+
+ // expected-error@+2{{'num' argument on 'worker' clause is not permitted on
a 'loop' construct associated with a 'parallel' compute construct}}
+#pragma acc parallel
+#pragma acc loop worker(num:i)
+ for(;;);
+
+#pragma acc serial
+#pragma acc loop worker
+ for(;;);
+
+ // expected-error@+2{{'num' argument on 'worker' clause is not permitted on
a 'loop' construct associated with a 'serial' compute construct}}
+#pragma acc serial
+#pragma acc loop worker(i)
+ for(;;);
+
+ // expected-error@+2{{'num' argument on 'worker' clause is not permitted on
a 'loop' construct associated with a 'serial' compute construct}}
+#pragma acc serial
+#pragma acc loop worker(num:i)
+ for(;;);
+
+#pragma acc kernels
+#pragma acc loop worker
+ for(;;);
+
+#pragma acc kernels
+#pragma acc loop worker(i)
+ for(;;);
+
+ Converts Cvts;
+
+#pragma acc kernels
+#pragma acc loop worker(Cvts)
+ for(;;);
+
+ NoConvert NoCvts;
+
+#pragma acc kernels
+ // expected-error@+1{{OpenACC clause 'worker' requires expression of integer
type ('NoConvert' invalid)}}
+#pragma acc loop worker(NoCvts)
+ for(;;);
+
+#pragma acc kernels
+#pragma acc loop worker(num:i)
+ for(;;);
+
+ // expected-error@+3{{'num' argument to 'worker' clause not allowed on a
'loop' construct associated with a 'kernels' construct that has a 'num_workers'
clause}}
+ // expected-note@+1{{previous clause is here}}
+#pragma acc kernels num_workers(i)
+#pragma acc loop worker(num:i)
+ for(;;);
+
+#pragma acc loop worker
+ for(;;) {
+ // expected-error@+3{{loop with a 'worker' clause may not exist in the
region of a 'worker' clause}}
+ // expected-error@+2{{loop with a 'gang' clause may not exist in the
region of a 'worker' clause}}
+ // expected-note@-4 2{{previous clause is here}}
+#pragma acc loop worker, gang
+ for(;;) {}
+ }
+
+#pragma acc loop worker
+ for(;;) {
+#pragma acc parallel
+#pragma acc loop worker, gang
+ for(;;) {}
+ }
+
+
+#pragma acc parallel
+#pragma acc loop worker
+ for(;;) {
+ // expected-error@+3{{loop with a 'worker' clause may not exist in the
region of a 'worker' clause}}
+ // expected-error@+2{{loop with a 'gang' clause may not exist in the
region of a 'worker' clause}}
+ // expected-note@-4 2{{previous clause is here}}
+#pragma acc loop worker, gang
+ for(;;) {}
+ }
+
+#pragma acc parallel
+#pragma acc loop worker
+ for(;;) {
+#pragma acc parallel
+#pragma acc loop worker, gang
+ for(;;) {}
+ }
+
+#pragma acc serial
+#pragma acc loop worker
+ for(;;) {
+ // expected-error@+3{{loop with a 'worker' clause may not exist in the
region of a 'worker' clause}}
+ // expected-error@+2{{loop with a 'gang' clause may not exist in the
region of a 'worker' clause}}
+ // expected-note@-4 2{{previous clause is here}}
+#pragma acc loop worker, gang
+ for(;;) {}
+ }
+
+#pragma acc serial
+#pragma acc loop worker
+ for(;;) {
+#pragma acc parallel
+#pragma acc loop worker, gang
+ for(;;) {}
+ }
+
+#pragma acc kernels
+#pragma acc loop worker
+ for(;;) {
+ // expected-error@+3{{loop with a 'worker' clause may not exist in the
region of a 'worker' clause}}
+ // expected-error@+2{{loop with a 'gang' clause may not exist in the
region of a 'worker' clause}}
+ // expected-note@-4 2{{previous clause is here}}
+#pragma acc loop worker, gang
+ for(;;) {}
+ }
+
+#pragma acc kernels
+#pragma acc loop worker
+ for(;;) {
+#pragma acc parallel
+#pragma acc loop worker, gang
+ for(;;) {}
+ }
+}
diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
index 2ffe47fbd74476..4461be86ea9996 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -2888,6 +2888,11 @@ void OpenACCClauseEnqueue::VisitAsyncClause(const
OpenACCAsyncClause &C) {
if (C.hasIntExpr())
Visitor.AddStmt(C.getIntExpr());
}
+
+void OpenACCClauseEnqueue::VisitWorkerClause(const OpenACCWorkerClause &C) {
+ if (C.hasIntExpr())
+ Visitor.AddStmt(C.getIntExpr());
+}
void OpenACCClauseEnqueue::VisitWaitClause(const OpenACCWaitClause &C) {
if (const Expr *DevNumExpr = C.getDevNumExpr())
Visitor.AddStmt(DevNumExpr);
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
