[clang] [OpenACC] Implement 'loop' 'vector' clause (PR #112259)
Erich Keane via cfe-commits
cfe-commits at lists.llvm.org
Mon Oct 14 13:49:55 PDT 2024
https://github.com/erichkeane created https://github.com/llvm/llvm-project/pull/112259
The 'vector' clause specifies the iterations to be executed in vector or SIMD mode. There are some limitations on which associated compute contexts may be associated with this and have arguments, but otherwise this is a fairly unrestricted clause.
It DOES have region limits like 'gang' and 'worker'.
>From 48e6540b70644bf37f331ed511a56bb7ec3f3aba Mon Sep 17 00:00:00 2001
From: erichkeane <ekeane at nvidia.com>
Date: Mon, 14 Oct 2024 09:12:20 -0700
Subject: [PATCH] [OpenACC] Implement 'loop' 'vector' clause
The 'vector' clause specifies the iterations to be executed in vector or
SIMD mode. There are some limitations on which associated compute
contexts may be associated with this and have arguments, but otherwise
this is a fairly unrestricted clause.
It DOES have region limits like 'gang' and 'worker'.
---
clang/include/clang/AST/OpenACCClause.h | 42 +--
.../clang/Basic/DiagnosticSemaKinds.td | 2 +-
clang/include/clang/Basic/OpenACCClauses.def | 1 +
clang/include/clang/Sema/SemaOpenACC.h | 6 +
clang/lib/AST/OpenACCClause.cpp | 31 +-
clang/lib/AST/StmtProfile.cpp | 6 +
clang/lib/AST/TextNodeDumper.cpp | 1 +
clang/lib/Sema/SemaOpenACC.cpp | 146 +++++++-
clang/lib/Sema/TreeTransform.h | 27 ++
clang/lib/Serialization/ASTReader.cpp | 7 +-
clang/lib/Serialization/ASTWriter.cpp | 9 +-
.../AST/ast-print-openacc-loop-construct.cpp | 75 ++++
clang/test/ParserOpenACC/parse-clauses.c | 11 +-
...p-construct-auto_seq_independent-clauses.c | 15 +-
.../loop-construct-device_type-clause.c | 1 -
.../SemaOpenACC/loop-construct-vector-ast.cpp | 346 ++++++++++++++++++
.../loop-construct-vector-clause.cpp | 136 +++++++
clang/tools/libclang/CIndex.cpp | 6 +
18 files changed, 806 insertions(+), 62 deletions(-)
create mode 100644 clang/test/SemaOpenACC/loop-construct-vector-ast.cpp
create mode 100644 clang/test/SemaOpenACC/loop-construct-vector-clause.cpp
diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h
index e8b8f477f91ae7..5ad4c336b6c531 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -119,32 +119,6 @@ class OpenACCSeqClause : public OpenACCClause {
}
};
-// Not yet implemented, but the type name is necessary for 'seq' diagnostics, so
-// this provides a basic, do-nothing implementation. We still need to add this
-// type to the visitors/etc, as well as get it to take its proper arguments.
-class OpenACCVectorClause : public OpenACCClause {
-protected:
- OpenACCVectorClause(SourceLocation BeginLoc, SourceLocation EndLoc)
- : OpenACCClause(OpenACCClauseKind::Vector, BeginLoc, EndLoc) {
- llvm_unreachable("Not yet implemented");
- }
-
-public:
- static bool classof(const OpenACCClause *C) {
- return C->getClauseKind() == OpenACCClauseKind::Vector;
- }
-
- static OpenACCVectorClause *
- 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 '('.
@@ -531,6 +505,22 @@ class OpenACCWorkerClause : public OpenACCClauseWithSingleIntExpr {
SourceLocation EndLoc);
};
+class OpenACCVectorClause : public OpenACCClauseWithSingleIntExpr {
+protected:
+ OpenACCVectorClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
+ Expr *IntExpr, SourceLocation EndLoc);
+
+public:
+ static bool classof(const OpenACCClause *C) {
+ return C->getClauseKind() == OpenACCClauseKind::Vector;
+ }
+
+ static OpenACCVectorClause *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 c709795e7b21d8..e78acc8dc8c57b 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12702,7 +12702,7 @@ def err_acc_gang_dim_value
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' "
+ "'%select{num_gangs|num_workers|vector_length}1' "
"clause">;
def err_acc_clause_in_clause_region
: Error<"loop with a '%0' clause may not exist in the region of a '%1' "
diff --git a/clang/include/clang/Basic/OpenACCClauses.def b/clang/include/clang/Basic/OpenACCClauses.def
index 4c0b56dc13e625..c65ebed751cf14 100644
--- a/clang/include/clang/Basic/OpenACCClauses.def
+++ b/clang/include/clang/Basic/OpenACCClauses.def
@@ -54,6 +54,7 @@ VISIT_CLAUSE(Reduction)
VISIT_CLAUSE(Self)
VISIT_CLAUSE(Seq)
VISIT_CLAUSE(Tile)
+VISIT_CLAUSE(Vector)
VISIT_CLAUSE(VectorLength)
VISIT_CLAUSE(Wait)
VISIT_CLAUSE(Worker)
diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h
index e253610a84b0bf..d6e56f85e2d54e 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -123,6 +123,11 @@ class SemaOpenACC : public SemaBase {
/// permits us to implement the restriction of no further 'gang' or 'worker'
/// clauses.
SourceLocation LoopWorkerClauseLoc;
+ /// If there is a current 'active' loop construct with a 'vector' 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', 'vector', or
+ /// 'worker' clauses.
+ SourceLocation LoopVectorClauseLoc;
// Redeclaration of the version in OpenACCClause.h.
using DeviceTypeArgument = std::pair<IdentifierInfo *, SourceLocation>;
@@ -679,6 +684,7 @@ class SemaOpenACC : public SemaBase {
OpenACCDirectiveKind DirKind;
SourceLocation OldLoopGangClauseOnKernelLoc;
SourceLocation OldLoopWorkerClauseLoc;
+ SourceLocation OldLoopVectorClauseLoc;
llvm::SmallVector<OpenACCLoopConstruct *> ParentlessLoopConstructs;
LoopInConstructRAII LoopRAII;
diff --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp
index 638252fd811f1d..1299e4f807ceb1 100644
--- a/clang/lib/AST/OpenACCClause.cpp
+++ b/clang/lib/AST/OpenACCClause.cpp
@@ -44,8 +44,8 @@ bool OpenACCClauseWithCondition::classof(const OpenACCClause *C) {
bool OpenACCClauseWithSingleIntExpr::classof(const OpenACCClause *C) {
return OpenACCNumWorkersClause::classof(C) ||
OpenACCVectorLengthClause::classof(C) ||
- OpenACCWorkerClause::classof(C) || OpenACCCollapseClause::classof(C) ||
- OpenACCAsyncClause::classof(C);
+ OpenACCVectorClause::classof(C) || OpenACCWorkerClause::classof(C) ||
+ OpenACCCollapseClause::classof(C) || OpenACCAsyncClause::classof(C);
}
OpenACCDefaultClause *OpenACCDefaultClause::Create(const ASTContext &C,
OpenACCDefaultClauseKind K,
@@ -424,11 +424,24 @@ OpenACCWorkerClause *OpenACCWorkerClause::Create(const ASTContext &C,
return new (Mem) OpenACCWorkerClause(BeginLoc, LParenLoc, IntExpr, EndLoc);
}
+OpenACCVectorClause::OpenACCVectorClause(SourceLocation BeginLoc,
+ SourceLocation LParenLoc,
+ Expr *IntExpr, SourceLocation EndLoc)
+ : OpenACCClauseWithSingleIntExpr(OpenACCClauseKind::Vector, BeginLoc,
+ LParenLoc, IntExpr, EndLoc) {
+ assert((!IntExpr || IntExpr->isInstantiationDependent() ||
+ IntExpr->getType()->isIntegerType()) &&
+ "Int expression type not scalar/dependent");
+}
+
OpenACCVectorClause *OpenACCVectorClause::Create(const ASTContext &C,
SourceLocation BeginLoc,
+ SourceLocation LParenLoc,
+ Expr *IntExpr,
SourceLocation EndLoc) {
- void *Mem = C.Allocate(sizeof(OpenACCVectorClause));
- return new (Mem) OpenACCVectorClause(BeginLoc, EndLoc);
+ void *Mem =
+ C.Allocate(sizeof(OpenACCVectorClause), alignof(OpenACCVectorClause));
+ return new (Mem) OpenACCVectorClause(BeginLoc, LParenLoc, IntExpr, EndLoc);
}
//===----------------------------------------------------------------------===//
@@ -662,3 +675,13 @@ void OpenACCClausePrinter::VisitWorkerClause(const OpenACCWorkerClause &C) {
OS << ")";
}
}
+
+void OpenACCClausePrinter::VisitVectorClause(const OpenACCVectorClause &C) {
+ OS << "vector";
+
+ if (C.hasIntExpr()) {
+ OS << "(length: ";
+ printExpr(C.getIntExpr());
+ OS << ")";
+ }
+}
diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index 25b1cbb8590869..01f9a30d06cd71 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -2635,6 +2635,12 @@ void OpenACCClauseProfiler::VisitWorkerClause(
Profiler.VisitStmt(Clause.getIntExpr());
}
+void OpenACCClauseProfiler::VisitVectorClause(
+ const OpenACCVectorClause &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 beccb0615f0e9c..01cfb1f63f7087 100644
--- a/clang/lib/AST/TextNodeDumper.cpp
+++ b/clang/lib/AST/TextNodeDumper.cpp
@@ -421,6 +421,7 @@ void TextNodeDumper::Visit(const OpenACCClause *C) {
case OpenACCClauseKind::Seq:
case OpenACCClauseKind::Tile:
case OpenACCClauseKind::Worker:
+ case OpenACCClauseKind::Vector:
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/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 1b24331cbd87ca..22aedbc70df8cc 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -389,6 +389,18 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
return false;
}
}
+ case OpenACCClauseKind::Vector: {
+ switch (DirectiveKind) {
+ case OpenACCDirectiveKind::Loop:
+ case OpenACCDirectiveKind::ParallelLoop:
+ case OpenACCDirectiveKind::SerialLoop:
+ case OpenACCDirectiveKind::KernelsLoop:
+ case OpenACCDirectiveKind::Routine:
+ return true;
+ default:
+ return false;
+ }
+ }
}
default:
@@ -512,14 +524,6 @@ class SemaOpenACCClauseVisitor {
OpenACCClause *Visit(SemaOpenACC::OpenACCParsedClause &Clause) {
switch (Clause.getClauseKind()) {
- case OpenACCClauseKind::Vector: {
- // TODO OpenACC: These are only implemented enough for the 'seq'
- // diagnostic, otherwise treats itself as unimplemented. When we
- // implement these, we can remove them from here.
- DiagIfSeqClause(Clause);
- return isNotImplemented();
- }
-
#define VISIT_CLAUSE(CLAUSE_NAME) \
case OpenACCClauseKind::CLAUSE_NAME: \
return Visit##CLAUSE_NAME##Clause(Clause);
@@ -1035,6 +1039,97 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitIndependentClause(
Clause.getEndLoc());
}
+OpenACCClause *SemaOpenACCClauseVisitor::VisitVectorClause(
+ 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:
+ case OpenACCDirectiveKind::Parallel:
+ // No restriction on when 'parallel' can contain an argument.
+ break;
+ case OpenACCDirectiveKind::Serial:
+ // GCC disallows this, and there is no real good reason for us to permit
+ // it, so disallow until we come up with a use case that makes sense.
+ SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_int_arg_invalid)
+ << OpenACCClauseKind::Vector << "num" << /*serial=*/3;
+ IntExpr = nullptr;
+ break;
+ case OpenACCDirectiveKind::Kernels: {
+ const auto *Itr =
+ llvm::find_if(SemaRef.getActiveComputeConstructInfo().Clauses,
+ llvm::IsaPred<OpenACCVectorLengthClause>);
+ if (Itr != SemaRef.getActiveComputeConstructInfo().Clauses.end()) {
+ SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict)
+ << OpenACCClauseKind::Vector << /*vector_length=*/2;
+ 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.2: When the parent compute construct is a kernels
+ // construct, the gang clause behaves as follows. ... The region of a loop
+ // with a gang clause may not contain another loop with a gang clause unless
+ // within a nested compute region.
+ if (SemaRef.LoopGangClauseOnKernelLoc.isValid()) {
+ // This handles the 'inner loop' diagnostic, but we cannot set that we're on
+ // one of these until we get to the end of the construct.
+ SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_in_clause_region)
+ << OpenACCClauseKind::Vector << OpenACCClauseKind::Gang
+ << /*skip kernels construct info*/ 0;
+ 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::Vector << OpenACCClauseKind::Worker
+ << /*skip kernels construct info*/ 0;
+ SemaRef.Diag(SemaRef.LoopWorkerClauseLoc,
+ diag::note_acc_previous_clause_here);
+ return nullptr;
+ }
+ // OpenACC 3.3 2.9.4: The region of a loop with a 'vector' clause may not
+ // contain a loop with a gang, worker, or vector clause unless within a nested
+ // compute region.
+ if (SemaRef.LoopVectorClauseLoc.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::Vector << OpenACCClauseKind::Vector
+ << /*skip kernels construct info*/ 0;
+ SemaRef.Diag(SemaRef.LoopVectorClauseLoc,
+ diag::note_acc_previous_clause_here);
+ return nullptr;
+ }
+
+ return OpenACCVectorClause::Create(Ctx, Clause.getBeginLoc(),
+ Clause.getLParenLoc(), IntExpr,
+ Clause.getEndLoc());
+}
+
OpenACCClause *SemaOpenACCClauseVisitor::VisitWorkerClause(
SemaOpenACC::OpenACCParsedClause &Clause) {
if (DiagIfSeqClause(Clause))
@@ -1099,6 +1194,20 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitWorkerClause(
return nullptr;
}
+ // OpenACC 3.3 2.9.4: The region of a loop with a 'vector' clause may not
+ // contain a loop with a gang, worker, or vector clause unless within a nested
+ // compute region.
+ if (SemaRef.LoopVectorClauseLoc.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::Vector
+ << /*skip kernels construct info*/ 0;
+ SemaRef.Diag(SemaRef.LoopVectorClauseLoc,
+ diag::note_acc_previous_clause_here);
+ return nullptr;
+ }
+
return OpenACCWorkerClause::Create(Ctx, Clause.getBeginLoc(),
Clause.getLParenLoc(), IntExpr,
Clause.getEndLoc());
@@ -1193,6 +1302,20 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitGangClause(
return nullptr;
}
+ // OpenACC 3.3 2.9.4: The region of a loop with a 'vector' clause may not
+ // contain a loop with a gang, worker, or vector clause unless within a nested
+ // compute region.
+ if (SemaRef.LoopVectorClauseLoc.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::Vector
+ << /*kernels construct info*/ 1;
+ SemaRef.Diag(SemaRef.LoopVectorClauseLoc,
+ diag::note_acc_previous_clause_here);
+ return nullptr;
+ }
+
return OpenACCGangClause::Create(Ctx, Clause.getBeginLoc(),
Clause.getLParenLoc(), GangKinds, IntExprs,
Clause.getEndLoc());
@@ -1313,6 +1436,7 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
: SemaRef(S), OldActiveComputeConstructInfo(S.ActiveComputeConstructInfo),
DirKind(DK), OldLoopGangClauseOnKernelLoc(S.LoopGangClauseOnKernelLoc),
OldLoopWorkerClauseLoc(S.LoopWorkerClauseLoc),
+ OldLoopVectorClauseLoc(S.LoopVectorClauseLoc),
LoopRAII(SemaRef, /*PreserveDepth=*/false) {
// Compute constructs end up taking their 'loop'.
if (DirKind == OpenACCDirectiveKind::Parallel ||
@@ -1330,6 +1454,7 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
// Implement the 'unless within a nested compute region' part.
SemaRef.LoopGangClauseOnKernelLoc = {};
SemaRef.LoopWorkerClauseLoc = {};
+ SemaRef.LoopVectorClauseLoc = {};
} else if (DirKind == OpenACCDirectiveKind::Loop) {
SetCollapseInfoBeforeAssociatedStmt(UnInstClauses, Clauses);
SetTileInfoBeforeAssociatedStmt(UnInstClauses, Clauses);
@@ -1355,6 +1480,10 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
auto *Itr = llvm::find_if(Clauses, llvm::IsaPred<OpenACCWorkerClause>);
if (Itr != Clauses.end())
SemaRef.LoopWorkerClauseLoc = (*Itr)->getBeginLoc();
+
+ auto *Itr2 = llvm::find_if(Clauses, llvm::IsaPred<OpenACCVectorClause>);
+ if (Itr2 != Clauses.end())
+ SemaRef.LoopVectorClauseLoc = (*Itr2)->getBeginLoc();
}
}
}
@@ -1429,6 +1558,7 @@ SemaOpenACC::AssociatedStmtRAII::~AssociatedStmtRAII() {
SemaRef.ActiveComputeConstructInfo = OldActiveComputeConstructInfo;
SemaRef.LoopGangClauseOnKernelLoc = OldLoopGangClauseOnKernelLoc;
SemaRef.LoopWorkerClauseLoc = OldLoopWorkerClauseLoc;
+ SemaRef.LoopVectorClauseLoc = OldLoopVectorClauseLoc;
if (DirKind == OpenACCDirectiveKind::Parallel ||
DirKind == OpenACCDirectiveKind::Serial ||
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index 45e8b3cf6bd8fc..c5a6e677ef8def 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -11815,6 +11815,33 @@ void OpenACCClauseTransform<Derived>::VisitWorkerClause(
ParsedClause.getEndLoc());
}
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitVectorClause(
+ const OpenACCVectorClause &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 = OpenACCVectorClause::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 ecc5d3c59a3549..1b2473f2457344 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -12345,10 +12345,15 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
return OpenACCWorkerClause::Create(getContext(), BeginLoc, LParenLoc,
WorkerExpr, EndLoc);
}
+ case OpenACCClauseKind::Vector: {
+ SourceLocation LParenLoc = readSourceLocation();
+ Expr *VectorExpr = readBool() ? readSubExpr() : nullptr;
+ return OpenACCVectorClause::Create(getContext(), BeginLoc, LParenLoc,
+ VectorExpr, EndLoc);
+ }
case OpenACCClauseKind::Finalize:
case OpenACCClauseKind::IfPresent:
- case OpenACCClauseKind::Vector:
case OpenACCClauseKind::NoHost:
case OpenACCClauseKind::UseDevice:
case OpenACCClauseKind::Delete:
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index 0a6e260e3e4e93..938d7b525cb959 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -8200,10 +8200,17 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
AddStmt(const_cast<Expr *>(WC->getIntExpr()));
return;
}
+ case OpenACCClauseKind::Vector: {
+ const auto *VC = cast<OpenACCVectorClause>(C);
+ writeSourceLocation(VC->getLParenLoc());
+ writeBool(VC->hasIntExpr());
+ if (VC->hasIntExpr())
+ AddStmt(const_cast<Expr *>(VC->getIntExpr()));
+ return;
+ }
case OpenACCClauseKind::Finalize:
case OpenACCClauseKind::IfPresent:
- case OpenACCClauseKind::Vector:
case OpenACCClauseKind::NoHost:
case OpenACCClauseKind::UseDevice:
case OpenACCClauseKind::Delete:
diff --git a/clang/test/AST/ast-print-openacc-loop-construct.cpp b/clang/test/AST/ast-print-openacc-loop-construct.cpp
index ee11435aaa4b1c..c0ca274f38dc2c 100644
--- a/clang/test/AST/ast-print-openacc-loop-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-loop-construct.cpp
@@ -216,4 +216,79 @@ void foo() {
#pragma acc kernels
#pragma acc loop worker(num:5)
for(;;);
+
+ // CHECK: #pragma acc loop vector
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc loop vector
+ for(;;);
+
+// CHECK: #pragma acc loop vector(length: 5)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc loop vector(5)
+ for(;;);
+
+// CHECK: #pragma acc loop vector(length: 5)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc loop vector(length:5)
+ for(;;);
+
+// CHECK: #pragma acc parallel
+// CHECK-NEXT: #pragma acc loop vector
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc parallel
+#pragma acc loop vector
+ for(;;);
+
+// CHECK: #pragma acc parallel
+// CHECK-NEXT: #pragma acc loop vector(length: 5)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc parallel
+#pragma acc loop vector(5)
+ for(;;);
+
+// CHECK: #pragma acc parallel
+// CHECK-NEXT: #pragma acc loop vector(length: 5)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc parallel
+#pragma acc loop vector(length:5)
+ for(;;);
+
+// CHECK: #pragma acc kernels
+// CHECK-NEXT: #pragma acc loop vector
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc kernels
+#pragma acc loop vector
+ for(;;);
+
+// CHECK: #pragma acc kernels
+// CHECK-NEXT: #pragma acc loop vector(length: 5)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc kernels
+#pragma acc loop vector(5)
+ for(;;);
+
+// CHECK: #pragma acc kernels
+// CHECK-NEXT: #pragma acc loop vector(length: 5)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc kernels
+#pragma acc loop vector(length:5)
+ for(;;);
+
+// CHECK: #pragma acc serial
+// CHECK-NEXT: #pragma acc loop vector
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc serial
+#pragma acc loop vector
+ for(;;);
+
}
diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c
index 81c48335cf0c42..26f39be80030b0 100644
--- a/clang/test/ParserOpenACC/parse-clauses.c
+++ b/clang/test/ParserOpenACC/parse-clauses.c
@@ -898,7 +898,6 @@ void IntExprParsing() {
#pragma acc set default_async(returns_int())
- // expected-warning at +1{{OpenACC clause 'vector' not yet implemented, clause ignored}}
#pragma acc loop vector
for(;;);
// expected-error at +1{{expected expression}}
@@ -908,8 +907,7 @@ void IntExprParsing() {
// expected-error at +1{{expected expression}}
#pragma acc loop vector(invalid:)
for(;;);
- // expected-error at +2{{invalid tag 'invalid' on 'vector' clause}}
- // expected-warning at +1{{OpenACC clause 'vector' not yet implemented, clause ignored}}
+ // expected-error at +1{{invalid tag 'invalid' on 'vector' clause}}
#pragma acc loop vector(invalid:5)
for(;;);
// expected-error at +1{{expected expression}}
@@ -932,20 +930,15 @@ void IntExprParsing() {
// expected-note at +1{{to match this '('}}
#pragma acc loop vector(num:6,4)
for(;;);
- // expected-warning at +1{{OpenACC clause 'vector' not yet implemented, clause ignored}}
#pragma acc loop vector(5)
for(;;);
- // expected-error at +2{{invalid tag 'num' on 'vector' clause}}
- // expected-warning at +1{{OpenACC clause 'vector' not yet implemented, clause ignored}}
+ // expected-error at +1{{invalid tag 'num' on 'vector' clause}}
#pragma acc loop vector(num:5)
for(;;);
- // expected-warning at +1{{OpenACC clause 'vector' not yet implemented, clause ignored}}
#pragma acc loop vector(length:5)
for(;;);
- // expected-warning at +1{{OpenACC clause 'vector' not yet implemented, clause ignored}}
#pragma acc loop vector(returns_int())
for(;;);
- // expected-warning at +1{{OpenACC clause 'vector' not yet implemented, clause ignored}}
#pragma acc loop vector(length:returns_int())
for(;;);
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 6a975956f3ff5c..ab10857e3cd858 100644
--- a/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
+++ b/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
@@ -45,7 +45,6 @@ void uses() {
for(;;);
#pragma acc loop auto worker
for(;;);
- // expected-warning at +1{{OpenACC clause 'vector' not yet implemented}}
#pragma acc loop auto vector
for(;;);
// expected-warning at +1{{OpenACC clause 'nohost' not yet implemented}}
@@ -181,7 +180,6 @@ void uses() {
for(;;);
#pragma acc loop worker auto
for(;;);
- // expected-warning at +1{{OpenACC clause 'vector' not yet implemented}}
#pragma acc loop vector auto
for(;;);
// expected-warning at +1{{OpenACC clause 'nohost' not yet implemented}}
@@ -318,7 +316,6 @@ void uses() {
for(;;);
#pragma acc loop independent worker
for(;;);
- // expected-warning at +1{{OpenACC clause 'vector' not yet implemented}}
#pragma acc loop independent vector
for(;;);
// expected-warning at +1{{OpenACC clause 'nohost' not yet implemented}}
@@ -454,7 +451,6 @@ void uses() {
for(;;);
#pragma acc loop worker independent
for(;;);
- // expected-warning at +1{{OpenACC clause 'vector' not yet implemented}}
#pragma acc loop vector independent
for(;;);
// expected-warning at +1{{OpenACC clause 'nohost' not yet implemented}}
@@ -591,9 +587,8 @@ void uses() {
// expected-note at +1{{previous clause is here}}
#pragma acc loop seq worker
for(;;);
- // expected-error at +3{{OpenACC clause 'vector' may not appear on the same construct as a 'seq' clause on a 'loop' construct}}
- // expected-note at +2{{previous clause is here}}
- // expected-warning at +1{{OpenACC clause 'vector' not yet implemented}}
+ // expected-error at +2{{OpenACC clause 'vector' may not appear on the same construct as a 'seq' clause on a 'loop' construct}}
+ // expected-note at +1{{previous clause is here}}
#pragma acc loop seq vector
for(;;);
// expected-warning at +1{{OpenACC clause 'finalize' not yet implemented}}
@@ -733,10 +728,8 @@ void uses() {
// expected-note at +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'.
- // TODOexpected-error at +3{{OpenACC clause 'vector' may not appear on the same construct as a 'seq' clause on a 'loop' construct}}
- // TODOexpected-note at +2{{previous clause is here}}
- // expected-warning at +1{{OpenACC clause 'vector' not yet implemented}}
+ // expected-error at +2{{OpenACC clause 'seq' may not appear on the same construct as a 'vector' clause on a 'loop' construct}}
+ // expected-note at +1{{previous clause is here}}
#pragma acc loop vector seq
for(;;);
// expected-warning at +1{{OpenACC clause 'finalize' not yet implemented}}
diff --git a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
index 51da8565f4e399..f60bf35a734fe8 100644
--- a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
@@ -38,7 +38,6 @@ void uses() {
// Only 'collapse', 'gang', 'worker', 'vector', 'seq', 'independent', 'auto',
// and 'tile' allowed after 'device_type'.
- // expected-warning at +1{{OpenACC clause 'vector' not yet implemented, clause ignored}}
#pragma acc loop device_type(*) vector
for(;;);
diff --git a/clang/test/SemaOpenACC/loop-construct-vector-ast.cpp b/clang/test/SemaOpenACC/loop-construct-vector-ast.cpp
new file mode 100644
index 00000000000000..390497cea0dd4e
--- /dev/null
+++ b/clang/test/SemaOpenACC/loop-construct-vector-ast.cpp
@@ -0,0 +1,346 @@
+// 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: vector clause
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+#pragma acc loop vector
+ for(;;);
+
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: vector 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 loop vector(I)
+ for(;;);
+
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: vector clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'ConvertsToInt' lvalue ParmVar{{.*}}'CTI' 'ConvertsToInt'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+#pragma acc loop vector(length:CTI)
+ for(;;);
+
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} parallel
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: vector clause
+ // 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
+#pragma acc parallel
+#pragma acc loop vector(length:IsI)
+ for(;;);
+
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} serial
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: vector 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 vector
+ for(;;);
+
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} kernels
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: vector clause
+ // 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
+#pragma acc kernels
+#pragma acc loop vector(length:IsI)
+ 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: vector clause
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+ //
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: vector 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
+ //
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: vector clause
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <UserDefinedConversion>
+ // CHECK-NEXT: CXXMemberCallExpr{{.*}}'int'
+ // CHECK-NEXT: MemberExpr{{.*}} .operator int
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Converts' lvalue ParmVar{{.*}}'CTI' 'Converts'
+ // 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: vector 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]+]]{{.*}} serial
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: vector 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: vector 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
+
+
+}
+
+struct Converts{
+ operator int();
+};
+
+void uses() {
+ // CHECK: FunctionDecl{{.*}} uses
+ // CHECK-NEXT: CompoundStmt
+
+ // CHECK-NEXT: CallExpr
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'void (*)(Converts, int)' <FunctionToPointerDecay>
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'void (Converts, int)' lvalue Function{{.*}} 'TemplUses' 'void (Converts, int)'
+ // CHECK-NEXT: CXXFunctionalCastExpr{{.*}} 'Converts' functional cast to Converts <NoOp>
+ // CHECK-NEXT: InitListExpr
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
+ TemplUses<3>(Converts{}, 5);
+
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}}
+ int i;
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}}
+ // CHECK-NEXT: CXXConstructExpr
+ Converts C;
+
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: vector clause
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+#pragma acc loop vector
+ for(;;);
+
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: vector 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 loop vector(i)
+ for(;;);
+
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: vector clause
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <UserDefinedConversion>
+ // CHECK-NEXT: CXXMemberCallExpr{{.*}}'int'
+ // CHECK-NEXT: MemberExpr{{.*}} .operator int
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Converts' lvalue Var{{.*}}'C' 'Converts'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+#pragma acc loop vector(length:C)
+ for(;;);
+
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} parallel
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: vector 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 vector
+ for(;;);
+
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} parallel
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: vector clause
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <UserDefinedConversion>
+ // CHECK-NEXT: CXXMemberCallExpr{{.*}}'int'
+ // CHECK-NEXT: MemberExpr{{.*}} .operator int
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Converts' lvalue Var{{.*}}'C' 'Converts'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+#pragma acc parallel
+#pragma acc loop vector(C)
+ for(;;);
+
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} parallel
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: vector 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 parallel
+#pragma acc loop vector(length:i)
+ for(;;);
+
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} kernels
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: vector clause
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+#pragma acc kernels
+#pragma acc loop vector
+ for(;;);
+
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} kernels
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: vector clause
+ // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <UserDefinedConversion>
+ // CHECK-NEXT: CXXMemberCallExpr{{.*}}'int'
+ // CHECK-NEXT: MemberExpr{{.*}} .operator int
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Converts' lvalue Var{{.*}}'C' 'Converts'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+#pragma acc kernels
+#pragma acc loop vector(C)
+ for(;;);
+
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} kernels
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: vector 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 vector(length:i)
+ for(;;);
+
+ // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} serial
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+ // CHECK-NEXT: vector 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 vector
+ for(;;);
+}
+#endif // PCH_HELPER
diff --git a/clang/test/SemaOpenACC/loop-construct-vector-clause.cpp b/clang/test/SemaOpenACC/loop-construct-vector-clause.cpp
new file mode 100644
index 00000000000000..ed2d4ee8ee0e34
--- /dev/null
+++ b/clang/test/SemaOpenACC/loop-construct-vector-clause.cpp
@@ -0,0 +1,136 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+template<typename Int, typename NotInt, typename ConvertsToInt>
+void TemplUses(Int I, NotInt NI, ConvertsToInt CTI) {
+#pragma acc loop vector(I)
+ for(;;);
+
+#pragma acc parallel
+#pragma acc loop vector(length: I)
+ for(;;);
+
+#pragma acc kernels
+#pragma acc loop vector(CTI)
+ for(;;);
+
+ // expected-error at +2{{OpenACC clause 'vector' requires expression of integer type ('NoConvert' invalid)}}
+#pragma acc kernels
+#pragma acc loop vector(length: NI)
+ for(;;);
+
+ // expected-error at +2{{'num' argument on 'vector' clause is not permitted on a 'loop' construct associated with a 'serial' compute construct}}
+#pragma acc serial
+#pragma acc loop vector(length: I)
+ for(;;);
+
+ // expected-error at +3{{'num' argument to 'vector' clause not allowed on a 'loop' construct associated with a 'kernels' construct that has a 'vector_length' clause}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels vector_length(I)
+#pragma acc loop vector(length: CTI)
+ for(;;);
+
+#pragma acc loop vector
+ for(;;) {
+ for(;;);
+ // expected-error at +2{{loop with a 'vector' clause may not exist in the region of a 'vector' clause}}
+ // expected-note at -4{{previous clause is here}}
+#pragma acc loop vector
+ for(;;);
+ for(;;);
+ }
+
+#pragma acc loop vector
+ for(;;) {
+ for(;;);
+ // expected-error at +4{{loop with a 'vector' clause may not exist in the region of a 'vector' clause}}
+ // expected-error at +3{{loop with a 'worker' clause may not exist in the region of a 'vector' clause}}
+ // expected-error at +2{{loop with a 'gang' clause may not exist in the region of a 'vector' clause}}
+ // expected-note at -6 3{{previous clause is here}}
+#pragma acc loop vector, worker, gang
+ for(;;);
+ for(;;);
+ }
+
+#pragma acc loop vector
+ for(;;) {
+#pragma acc serial
+#pragma acc loop vector
+ for(;;);
+ }
+}
+
+struct NoConvert{};
+struct Converts{
+ operator int();
+};
+
+void uses() {
+ TemplUses(5, NoConvert{}, Converts{}); // expected-note{{in instantiation of function template specialization}}
+
+ unsigned i;
+ NoConvert NI;
+ Converts CTI;
+
+#pragma acc loop vector(i)
+ for(;;);
+
+#pragma acc parallel
+#pragma acc loop vector(length: i)
+ for(;;);
+
+#pragma acc kernels
+#pragma acc loop vector(CTI)
+ for(;;);
+
+ // expected-error at +2{{OpenACC clause 'vector' requires expression of integer type ('NoConvert' invalid)}}
+#pragma acc kernels
+#pragma acc loop vector(length: NI)
+ for(;;);
+
+ // expected-error at +2{{'num' argument on 'vector' clause is not permitted on a 'loop' construct associated with a 'serial' compute construct}}
+#pragma acc serial
+#pragma acc loop vector(length: i)
+ for(;;);
+
+ // expected-error at +3{{'num' argument to 'vector' clause not allowed on a 'loop' construct associated with a 'kernels' construct that has a 'vector_length' clause}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels vector_length(i)
+#pragma acc loop vector(length: i)
+ for(;;);
+
+#pragma acc loop vector
+ for(;;) {
+ for(;;);
+ // expected-error at +2{{loop with a 'vector' clause may not exist in the region of a 'vector' clause}}
+ // expected-note at -4{{previous clause is here}}
+#pragma acc loop vector
+ for(;;);
+ for(;;);
+ }
+
+#pragma acc loop vector
+ for(;;) {
+#pragma acc serial
+#pragma acc loop vector
+ for(;;);
+ }
+
+#pragma acc loop vector
+ for(;;) {
+ for(;;);
+ // expected-error at +4{{loop with a 'vector' clause may not exist in the region of a 'vector' clause}}
+ // expected-error at +3{{loop with a 'worker' clause may not exist in the region of a 'vector' clause}}
+ // expected-error at +2{{loop with a 'gang' clause may not exist in the region of a 'vector' clause}}
+ // expected-note at -6 3{{previous clause is here}}
+#pragma acc loop vector, worker, gang
+ for(;;);
+ for(;;);
+ }
+
+#pragma acc loop vector
+ for(;;) {
+#pragma acc serial
+#pragma acc loop vector, worker, gang
+ for(;;);
+ }
+}
diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
index 4461be86ea9996..55dd3d7b691143 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -2893,6 +2893,12 @@ void OpenACCClauseEnqueue::VisitWorkerClause(const OpenACCWorkerClause &C) {
if (C.hasIntExpr())
Visitor.AddStmt(C.getIntExpr());
}
+
+void OpenACCClauseEnqueue::VisitVectorClause(const OpenACCVectorClause &C) {
+ if (C.hasIntExpr())
+ Visitor.AddStmt(C.getIntExpr());
+}
+
void OpenACCClauseEnqueue::VisitWaitClause(const OpenACCWaitClause &C) {
if (const Expr *DevNumExpr = C.getDevNumExpr())
Visitor.AddStmt(DevNumExpr);
More information about the cfe-commits
mailing list