[clang] a13c514 - [OpenACC] Implement firstprivate clause for compute constructs
via cfe-commits
cfe-commits at lists.llvm.org
Fri May 3 06:33:42 PDT 2024
Author: erichkeane
Date: 2024-05-03T06:33:35-07:00
New Revision: a13c5140a2a26923f3e7bf3684409425ff54de6f
URL: https://github.com/llvm/llvm-project/commit/a13c5140a2a26923f3e7bf3684409425ff54de6f
DIFF: https://github.com/llvm/llvm-project/commit/a13c5140a2a26923f3e7bf3684409425ff54de6f.diff
LOG: [OpenACC] Implement firstprivate clause for compute constructs
This clause is pretty nearly copy/paste from private, except that it
doesn't support 'loop', and thus 'kernelsloop' for appertainment.
Added:
clang/test/SemaOpenACC/compute-construct-firstprivate-clause.c
clang/test/SemaOpenACC/compute-construct-firstprivate-clause.cpp
Modified:
clang/include/clang/AST/OpenACCClause.h
clang/include/clang/Basic/OpenACCClauses.def
clang/include/clang/Sema/SemaOpenACC.h
clang/lib/AST/OpenACCClause.cpp
clang/lib/AST/StmtProfile.cpp
clang/lib/AST/TextNodeDumper.cpp
clang/lib/Parse/ParseOpenACC.cpp
clang/lib/Sema/SemaOpenACC.cpp
clang/lib/Sema/TreeTransform.h
clang/lib/Serialization/ASTReader.cpp
clang/lib/Serialization/ASTWriter.cpp
clang/test/AST/ast-print-openacc-compute-construct.cpp
clang/test/ParserOpenACC/parse-clauses.c
clang/test/SemaOpenACC/compute-construct-varlist-ast.cpp
clang/tools/libclang/CIndex.cpp
Removed:
################################################################################
diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h
index bb4cb1f5d5080b..df290d06efd203 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -294,6 +294,25 @@ class OpenACCPrivateClause final
ArrayRef<Expr *> VarList, SourceLocation EndLoc);
};
+class OpenACCFirstPrivateClause final
+ : public OpenACCClauseWithVarList,
+ public llvm::TrailingObjects<OpenACCFirstPrivateClause, Expr *> {
+
+ OpenACCFirstPrivateClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
+ ArrayRef<Expr *> VarList, SourceLocation EndLoc)
+ : OpenACCClauseWithVarList(OpenACCClauseKind::FirstPrivate, BeginLoc,
+ LParenLoc, EndLoc) {
+ std::uninitialized_copy(VarList.begin(), VarList.end(),
+ getTrailingObjects<Expr *>());
+ setExprs(MutableArrayRef(getTrailingObjects<Expr *>(), VarList.size()));
+ }
+
+public:
+ static OpenACCFirstPrivateClause *
+ Create(const ASTContext &C, SourceLocation BeginLoc, SourceLocation LParenLoc,
+ ArrayRef<Expr *> VarList, SourceLocation EndLoc);
+};
+
template <class Impl> class OpenACCClauseVisitor {
Impl &getDerived() { return static_cast<Impl &>(*this); }
diff --git a/clang/include/clang/Basic/OpenACCClauses.def b/clang/include/clang/Basic/OpenACCClauses.def
index 6c3c2db66ef0cf..884f0297aa50e0 100644
--- a/clang/include/clang/Basic/OpenACCClauses.def
+++ b/clang/include/clang/Basic/OpenACCClauses.def
@@ -16,11 +16,12 @@
// VISIT_CLAUSE(CLAUSE_NAME)
VISIT_CLAUSE(Default)
+VISIT_CLAUSE(FirstPrivate)
VISIT_CLAUSE(If)
-VISIT_CLAUSE(Self)
VISIT_CLAUSE(NumGangs)
VISIT_CLAUSE(NumWorkers)
VISIT_CLAUSE(Private)
+VISIT_CLAUSE(Self)
VISIT_CLAUSE(VectorLength)
#undef VISIT_CLAUSE
diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h
index edb0cbb7c5d552..a25c5244d9607a 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -117,7 +117,8 @@ class SemaOpenACC : public SemaBase {
}
ArrayRef<Expr *> getVarList() {
- assert(ClauseKind == OpenACCClauseKind::Private &&
+ assert((ClauseKind == OpenACCClauseKind::Private ||
+ ClauseKind == OpenACCClauseKind::FirstPrivate) &&
"Parsed clause kind does not have a var-list");
return std::get<VarListDetails>(Details).VarList;
}
@@ -165,13 +166,15 @@ class SemaOpenACC : public SemaBase {
}
void setVarListDetails(ArrayRef<Expr *> VarList) {
- assert(ClauseKind == OpenACCClauseKind::Private &&
+ assert((ClauseKind == OpenACCClauseKind::Private ||
+ ClauseKind == OpenACCClauseKind::FirstPrivate) &&
"Parsed clause kind does not have a var-list");
Details = VarListDetails{{VarList.begin(), VarList.end()}};
}
void setVarListDetails(llvm::SmallVector<Expr *> &&VarList) {
- assert(ClauseKind == OpenACCClauseKind::Private &&
+ assert((ClauseKind == OpenACCClauseKind::Private ||
+ ClauseKind == OpenACCClauseKind::FirstPrivate) &&
"Parsed clause kind does not have a var-list");
Details = VarListDetails{std::move(VarList)};
}
diff --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp
index 885f3b7618ec55..fdb802c8d2a66f 100644
--- a/clang/lib/AST/OpenACCClause.cpp
+++ b/clang/lib/AST/OpenACCClause.cpp
@@ -144,6 +144,15 @@ OpenACCPrivateClause *OpenACCPrivateClause::Create(const ASTContext &C,
return new (Mem) OpenACCPrivateClause(BeginLoc, LParenLoc, VarList, EndLoc);
}
+OpenACCFirstPrivateClause *OpenACCFirstPrivateClause::Create(
+ const ASTContext &C, SourceLocation BeginLoc, SourceLocation LParenLoc,
+ ArrayRef<Expr *> VarList, SourceLocation EndLoc) {
+ void *Mem = C.Allocate(
+ OpenACCFirstPrivateClause::totalSizeToAlloc<Expr *>(VarList.size()));
+ return new (Mem)
+ OpenACCFirstPrivateClause(BeginLoc, LParenLoc, VarList, EndLoc);
+}
+
//===----------------------------------------------------------------------===//
// OpenACC clauses printing methods
//===----------------------------------------------------------------------===//
@@ -198,3 +207,11 @@ void OpenACCClausePrinter::VisitPrivateClause(const OpenACCPrivateClause &C) {
[&](const Expr *E) { printExpr(E); });
OS << ")";
}
+
+void OpenACCClausePrinter::VisitFirstPrivateClause(
+ const OpenACCFirstPrivateClause &C) {
+ OS << "firstprivate(";
+ llvm::interleaveComma(C.getVarList(), OS,
+ [&](const Expr *E) { printExpr(E); });
+ OS << ")";
+}
diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index 973f6f97bae0bf..352bef47b39e45 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -2515,6 +2515,12 @@ void OpenACCClauseProfiler::VisitPrivateClause(
Profiler.VisitStmt(E);
}
+void OpenACCClauseProfiler::VisitFirstPrivateClause(
+ const OpenACCFirstPrivateClause &Clause) {
+ for (auto *E : Clause.getVarList())
+ Profiler.VisitStmt(E);
+}
+
void OpenACCClauseProfiler::VisitVectorLengthClause(
const OpenACCVectorLengthClause &Clause) {
assert(Clause.hasIntExpr() &&
diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp
index 89f50d6dacfd23..a2cd5040230060 100644
--- a/clang/lib/AST/TextNodeDumper.cpp
+++ b/clang/lib/AST/TextNodeDumper.cpp
@@ -398,10 +398,11 @@ void TextNodeDumper::Visit(const OpenACCClause *C) {
OS << '(' << cast<OpenACCDefaultClause>(C)->getDefaultClauseKind() << ')';
break;
case OpenACCClauseKind::If:
- case OpenACCClauseKind::Self:
+ case OpenACCClauseKind::FirstPrivate:
case OpenACCClauseKind::NumGangs:
case OpenACCClauseKind::NumWorkers:
case OpenACCClauseKind::Private:
+ case OpenACCClauseKind::Self:
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 2d1ec6539b2fd1..8784acc9a2f37f 100644
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@@ -926,7 +926,6 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
case OpenACCClauseKind::Device:
case OpenACCClauseKind::DeviceResident:
case OpenACCClauseKind::DevicePtr:
- case OpenACCClauseKind::FirstPrivate:
case OpenACCClauseKind::Host:
case OpenACCClauseKind::Link:
case OpenACCClauseKind::NoCreate:
@@ -934,6 +933,7 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
case OpenACCClauseKind::UseDevice:
ParseOpenACCVarList();
break;
+ case OpenACCClauseKind::FirstPrivate:
case OpenACCClauseKind::Private:
ParsedClause.setVarListDetails(ParseOpenACCVarList());
break;
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index e07c5b2317f279..778e09e7ce0080 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -104,6 +104,16 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
default:
return false;
}
+ case OpenACCClauseKind::FirstPrivate:
+ switch (DirectiveKind) {
+ case OpenACCDirectiveKind::Parallel:
+ case OpenACCDirectiveKind::Serial:
+ case OpenACCDirectiveKind::ParallelLoop:
+ case OpenACCDirectiveKind::SerialLoop:
+ return true;
+ default:
+ return false;
+ }
case OpenACCClauseKind::Private:
switch (DirectiveKind) {
case OpenACCDirectiveKind::Parallel:
@@ -331,6 +341,21 @@ SemaOpenACC::ActOnClause(ArrayRef<const OpenACCClause *> ExistingClauses,
getASTContext(), Clause.getBeginLoc(), Clause.getLParenLoc(),
Clause.getVarList(), Clause.getEndLoc());
}
+ case OpenACCClauseKind::FirstPrivate: {
+ // Restrictions only properly implemented on 'compute' constructs, and
+ // 'compute' constructs are the only construct that can do anything with
+ // this yet, so skip/treat as unimplemented in this case.
+ if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()))
+ break;
+
+ // ActOnVar ensured that everything is a valid variable reference, so there
+ // really isn't anything to do here. GCC does some duplicate-finding, though
+ // it isn't apparent in the standard where this is justified.
+
+ return OpenACCFirstPrivateClause::Create(
+ getASTContext(), Clause.getBeginLoc(), Clause.getLParenLoc(),
+ Clause.getVarList(), Clause.getEndLoc());
+ }
default:
break;
}
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index dff7e9df636b96..6ed5d62b7dbd7b 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -11112,6 +11112,23 @@ class OpenACCClauseTransform final
SemaOpenACC::OpenACCParsedClause &ParsedClause;
OpenACCClause *NewClause = nullptr;
+ llvm::SmallVector<Expr *> VisitVarList(ArrayRef<Expr *> VarList) {
+ llvm::SmallVector<Expr *> InstantiatedVarList;
+ for (Expr *CurVar : VarList) {
+ ExprResult Res = Self.TransformExpr(CurVar);
+
+ if (!Res.isUsable())
+ continue;
+
+ Res = Self.getSema().OpenACC().ActOnVar(Res.get());
+
+ if (Res.isUsable())
+ InstantiatedVarList.push_back(Res.get());
+ }
+
+ return InstantiatedVarList;
+ }
+
public:
OpenACCClauseTransform(TreeTransform<Derived> &Self,
ArrayRef<const OpenACCClause *> ExistingClauses,
@@ -11206,22 +11223,20 @@ void OpenACCClauseTransform<Derived>::VisitNumGangsClause(
template <typename Derived>
void OpenACCClauseTransform<Derived>::VisitPrivateClause(
const OpenACCPrivateClause &C) {
- llvm::SmallVector<Expr *> InstantiatedVarList;
-
- for (Expr *CurVar : C.getVarList()) {
- ExprResult Res = Self.TransformExpr(CurVar);
-
- if (!Res.isUsable())
- return;
+ ParsedClause.setVarListDetails(VisitVarList(C.getVarList()));
- Res = Self.getSema().OpenACC().ActOnVar(Res.get());
+ NewClause = OpenACCPrivateClause::Create(
+ Self.getSema().getASTContext(), ParsedClause.getBeginLoc(),
+ ParsedClause.getLParenLoc(), ParsedClause.getVarList(),
+ ParsedClause.getEndLoc());
+}
- if (Res.isUsable())
- InstantiatedVarList.push_back(Res.get());
- }
- ParsedClause.setVarListDetails(std::move(InstantiatedVarList));
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitFirstPrivateClause(
+ const OpenACCFirstPrivateClause &C) {
+ ParsedClause.setVarListDetails(VisitVarList(C.getVarList()));
- NewClause = OpenACCPrivateClause::Create(
+ NewClause = OpenACCFirstPrivateClause::Create(
Self.getSema().getASTContext(), ParsedClause.getBeginLoc(),
ParsedClause.getLParenLoc(), ParsedClause.getVarList(),
ParsedClause.getEndLoc());
diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
index 29b81c1a753cf5..e80ea77f568d12 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -11835,6 +11835,12 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
return OpenACCPrivateClause::Create(getContext(), BeginLoc, LParenLoc,
VarList, EndLoc);
}
+ case OpenACCClauseKind::FirstPrivate: {
+ SourceLocation LParenLoc = readSourceLocation();
+ llvm::SmallVector<Expr *> VarList = readOpenACCVarList();
+ return OpenACCFirstPrivateClause::Create(getContext(), BeginLoc, LParenLoc,
+ VarList, EndLoc);
+ }
case OpenACCClauseKind::Finalize:
case OpenACCClauseKind::IfPresent:
case OpenACCClauseKind::Seq:
@@ -11851,7 +11857,6 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
case OpenACCClauseKind::Device:
case OpenACCClauseKind::DevicePtr:
case OpenACCClauseKind::DeviceResident:
- case OpenACCClauseKind::FirstPrivate:
case OpenACCClauseKind::Host:
case OpenACCClauseKind::Link:
case OpenACCClauseKind::NoCreate:
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index 80c7ce643088b6..df53a1e4f2e1b7 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -7787,6 +7787,12 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
writeOpenACCVarList(PC);
return;
}
+ case OpenACCClauseKind::FirstPrivate: {
+ const auto *FPC = cast<OpenACCFirstPrivateClause>(C);
+ writeSourceLocation(FPC->getLParenLoc());
+ writeOpenACCVarList(FPC);
+ return;
+ }
case OpenACCClauseKind::Finalize:
case OpenACCClauseKind::IfPresent:
case OpenACCClauseKind::Seq:
@@ -7803,7 +7809,6 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
case OpenACCClauseKind::Device:
case OpenACCClauseKind::DevicePtr:
case OpenACCClauseKind::DeviceResident:
- case OpenACCClauseKind::FirstPrivate:
case OpenACCClauseKind::Host:
case OpenACCClauseKind::Link:
case OpenACCClauseKind::NoCreate:
diff --git a/clang/test/AST/ast-print-openacc-compute-construct.cpp b/clang/test/AST/ast-print-openacc-compute-construct.cpp
index cd39ea087b3cad..993f4950217b5e 100644
--- a/clang/test/AST/ast-print-openacc-compute-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-compute-construct.cpp
@@ -38,5 +38,9 @@ void foo() {
// CHECK: #pragma acc parallel private(i, array[1], array, array[1:2])
#pragma acc parallel private(i, array[1], array, array[1:2])
while(true);
+
+// CHECK: #pragma acc parallel firstprivate(i, array[1], array, array[1:2])
+#pragma acc parallel firstprivate(i, array[1], array, array[1:2])
+ while(true);
}
diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c
index fa43f42585fc2b..e67031fd0cb222 100644
--- a/clang/test/ParserOpenACC/parse-clauses.c
+++ b/clang/test/ParserOpenACC/parse-clauses.c
@@ -598,13 +598,11 @@ void VarListClauses() {
#pragma acc serial private(s.array[s.value : 5], s.value), seq
for(;;){}
- // expected-error at +3{{expected ','}}
- // expected-warning at +2{{OpenACC clause 'firstprivate' not yet implemented, clause ignored}}
+ // expected-error at +2{{expected ','}}
// expected-warning at +1{{OpenACC clause 'seq' not yet implemented, clause ignored}}
#pragma acc serial firstprivate(s.array[s.value] s.array[s.value :5] ), seq
for(;;){}
- // expected-warning at +2{{OpenACC clause 'firstprivate' not yet implemented, clause ignored}}
// expected-warning at +1{{OpenACC clause 'seq' not yet implemented, clause ignored}}
#pragma acc serial firstprivate(s.array[s.value : 5], s.value), seq
for(;;){}
diff --git a/clang/test/SemaOpenACC/compute-construct-firstprivate-clause.c b/clang/test/SemaOpenACC/compute-construct-firstprivate-clause.c
new file mode 100644
index 00000000000000..4e057bf32c2d6d
--- /dev/null
+++ b/clang/test/SemaOpenACC/compute-construct-firstprivate-clause.c
@@ -0,0 +1,55 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+typedef struct IsComplete {
+ struct S { int A; } CompositeMember;
+ int ScalarMember;
+ float ArrayMember[5];
+ void *PointerMember;
+} Complete;
+void uses(int IntParam, short *PointerParam, float ArrayParam[5], Complete CompositeParam) {
+ int LocalInt;
+ short *LocalPointer;
+ float LocalArray[5];
+ Complete LocalComposite;
+ // Check Appertainment:
+#pragma acc parallel firstprivate(LocalInt)
+ while(1);
+#pragma acc serial firstprivate(LocalInt)
+ while(1);
+ // expected-error at +1{{OpenACC 'firstprivate' clause is not valid on 'kernels' directive}}
+#pragma acc kernels firstprivate(LocalInt)
+ while(1);
+
+ // Valid cases:
+#pragma acc parallel firstprivate(LocalInt, LocalPointer, LocalArray)
+ while(1);
+#pragma acc parallel firstprivate(LocalArray[2:1])
+ while(1);
+
+#pragma acc parallel firstprivate(LocalComposite.ScalarMember, LocalComposite.ScalarMember)
+ while(1);
+
+ // expected-error at +1{{OpenACC variable is not a valid variable name, sub-array, array element, or composite variable member}}
+#pragma acc parallel firstprivate(1 + IntParam)
+ while(1);
+
+ // expected-error at +1{{OpenACC variable is not a valid variable name, sub-array, array element, or composite variable member}}
+#pragma acc parallel firstprivate(+IntParam)
+ while(1);
+
+ // expected-error at +1{{OpenACC sub-array length is unspecified and cannot be inferred because the subscripted value is not an array}}
+#pragma acc parallel firstprivate(PointerParam[2:])
+ while(1);
+
+ // expected-error at +1{{OpenACC sub-array specified range [2:5] would be out of the range of the subscripted array size of 5}}
+#pragma acc parallel firstprivate(ArrayParam[2:5])
+ while(1);
+
+ // expected-error at +2{{OpenACC sub-array specified range [2:5] would be out of the range of the subscripted array size of 5}}
+ // expected-error at +1{{OpenACC variable is not a valid variable name, sub-array, array element, or composite variable member}}
+#pragma acc parallel firstprivate((float*)ArrayParam[2:5])
+ while(1);
+ // expected-error at +1{{OpenACC variable is not a valid variable name, sub-array, array element, or composite variable member}}
+#pragma acc parallel firstprivate((float)ArrayParam[2])
+ while(1);
+}
diff --git a/clang/test/SemaOpenACC/compute-construct-firstprivate-clause.cpp b/clang/test/SemaOpenACC/compute-construct-firstprivate-clause.cpp
new file mode 100644
index 00000000000000..2fbb80f7b2fbd6
--- /dev/null
+++ b/clang/test/SemaOpenACC/compute-construct-firstprivate-clause.cpp
@@ -0,0 +1,113 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+enum SomeE{};
+typedef struct IsComplete {
+ struct S { int A; } CompositeMember;
+ int ScalarMember;
+ float ArrayMember[5];
+ SomeE EnumMember;
+ char *PointerMember;
+} Complete;
+
+void uses(int IntParam, char *PointerParam, float ArrayParam[5], Complete CompositeParam, int &IntParamRef) {
+ int LocalInt;
+ char *LocalPointer;
+ float LocalArray[5];
+ // Check Appertainment:
+#pragma acc parallel firstprivate(LocalInt)
+ while(1);
+#pragma acc serial firstprivate(LocalInt)
+ while(1);
+ // expected-error at +1{{OpenACC 'firstprivate' clause is not valid on 'kernels' directive}}
+#pragma acc kernels firstprivate(LocalInt)
+ while(1);
+
+ // Valid cases:
+#pragma acc parallel firstprivate(LocalInt, LocalPointer, LocalArray)
+ while(1);
+#pragma acc parallel firstprivate(LocalArray[2:1])
+ while(1);
+
+ Complete LocalComposite2;
+#pragma acc parallel firstprivate(LocalComposite2.ScalarMember, LocalComposite2.ScalarMember)
+ while(1);
+
+ // expected-error at +1{{OpenACC variable is not a valid variable name, sub-array, array element, or composite variable member}}
+#pragma acc parallel firstprivate(1 + IntParam)
+ while(1);
+
+ // expected-error at +1{{OpenACC variable is not a valid variable name, sub-array, array element, or composite variable member}}
+#pragma acc parallel firstprivate(+IntParam)
+ while(1);
+
+ // expected-error at +1{{OpenACC sub-array length is unspecified and cannot be inferred because the subscripted value is not an array}}
+#pragma acc parallel firstprivate(PointerParam[2:])
+ while(1);
+
+ // expected-error at +1{{OpenACC sub-array specified range [2:5] would be out of the range of the subscripted array size of 5}}
+#pragma acc parallel firstprivate(ArrayParam[2:5])
+ while(1);
+
+ // expected-error at +2{{OpenACC sub-array specified range [2:5] would be out of the range of the subscripted array size of 5}}
+ // expected-error at +1{{OpenACC variable is not a valid variable name, sub-array, array element, or composite variable member}}
+#pragma acc parallel firstprivate((float*)ArrayParam[2:5])
+ while(1);
+ // expected-error at +1{{OpenACC variable is not a valid variable name, sub-array, array element, or composite variable member}}
+#pragma acc parallel firstprivate((float)ArrayParam[2])
+ while(1);
+}
+
+template<typename T, unsigned I, typename V>
+void TemplUses(T t, T (&arrayT)[I], V TemplComp) {
+ // expected-error at +1{{OpenACC variable is not a valid variable name, sub-array, array element, or composite variable member}}
+#pragma acc parallel private(+t)
+ while(true);
+
+ // NTTP's are only valid if it is a reference to something.
+ // expected-error at +2{{OpenACC variable is not a valid variable name, sub-array, array element, or composite variable member}}
+ // expected-note@#TEMPL_USES_INST{{in instantiation of}}
+#pragma acc parallel private(I)
+ while(true);
+
+ // expected-error at +1{{OpenACC variable is not a valid variable name, sub-array, array element, or composite variable member}}
+#pragma acc parallel private(t, I)
+ while(true);
+
+#pragma acc parallel private(arrayT)
+ while(true);
+
+#pragma acc parallel private(TemplComp)
+ while(true);
+
+#pragma acc parallel private(TemplComp.PointerMember[5])
+ while(true);
+ int *Pointer;
+#pragma acc parallel private(Pointer[:I])
+ while(true);
+#pragma acc parallel private(Pointer[:t])
+ while(true);
+ // expected-error at +1{{OpenACC sub-array length is unspecified and cannot be inferred because the subscripted value is not an array}}
+#pragma acc parallel private(Pointer[1:])
+ while(true);
+}
+
+template<unsigned I, auto &NTTP_REF>
+void NTTP() {
+ // NTTP's are only valid if it is a reference to something.
+ // expected-error at +2{{OpenACC variable is not a valid variable name, sub-array, array element, or composite variable member}}
+ // expected-note@#NTTP_INST{{in instantiation of}}
+#pragma acc parallel private(I)
+ while(true);
+
+#pragma acc parallel private(NTTP_REF)
+ while(true);
+}
+
+void Inst() {
+ static constexpr int NTTP_REFed = 1;
+ int i;
+ int Arr[5];
+ Complete C;
+ TemplUses(i, Arr, C); // #TEMPL_USES_INST
+ NTTP<5, NTTP_REFed>(); // #NTTP_INST
+}
diff --git a/clang/test/SemaOpenACC/compute-construct-varlist-ast.cpp b/clang/test/SemaOpenACC/compute-construct-varlist-ast.cpp
index aad4b9963c70db..8d9b678c8284fb 100644
--- a/clang/test/SemaOpenACC/compute-construct-varlist-ast.cpp
+++ b/clang/test/SemaOpenACC/compute-construct-varlist-ast.cpp
@@ -35,6 +35,20 @@ void NormalUses(float *PointerParam) {
// CHECK-NEXT: CXXBoolLiteralExpr
// CHECK-NEXT: NullStmt
+#pragma acc parallel firstprivate(GlobalArray, PointerParam[Global])
+ while(true);
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+ // CHECK-NEXT: firstprivate clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'short[5]' lvalue Var{{.*}}'GlobalArray' 'short[5]'
+ // CHECK-NEXT: ArraySubscriptExpr{{.*}}'float' lvalue
+ // CHECK-NEXT: ImplicitCastExpr{{.*}} 'float *' <LValueToRValue>
+ // CHECK-NEXT: DeclRefExpr{{.*}}'float *' lvalue ParmVar{{.*}}'PointerParam' 'float *'
+ // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
+ // CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue Var{{.*}}'Global' 'int'
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: NullStmt
+
#pragma acc parallel private(GlobalArray) private(PointerParam[Global])
while(true);
// CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
@@ -65,6 +79,22 @@ void NormalUses(float *PointerParam) {
// CHECK-NEXT: WhileStmt
// CHECK-NEXT: CXXBoolLiteralExpr
// CHECK-NEXT: NullStmt
+
+#pragma acc parallel firstprivate(GlobalArray, PointerParam[Global : Global])
+ while(true);
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+ // CHECK-NEXT: firstprivate clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'short[5]' lvalue Var{{.*}}'GlobalArray' 'short[5]'
+ // CHECK-NEXT: ArraySectionExpr
+ // CHECK-NEXT: ImplicitCastExpr{{.*}} 'float *' <LValueToRValue>
+ // CHECK-NEXT: DeclRefExpr{{.*}}'float *' lvalue ParmVar{{.*}} 'PointerParam' 'float *'
+ // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
+ // CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue Var{{.*}}'Global' 'int'
+ // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
+ // CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue Var{{.*}}'Global' 'int'
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: NullStmt
}
// This example is an error typically, but we want to make sure we're properly
@@ -83,6 +113,14 @@ void UnInstTempl() {
// CHECK-NEXT: WhileStmt
// CHECK-NEXT: CXXBoolLiteralExpr
// CHECK-NEXT: NullStmt
+#pragma acc parallel firstprivate(I)
+ while(true);
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+ // CHECK-NEXT: firstprivate clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'unsigned int' NonTypeTemplateParm{{.*}}'I' 'unsigned int'
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: NullStmt
}
template<auto &NTTP, typename T, typename U>
@@ -120,6 +158,16 @@ void TemplUses(T t, U u, T*PointerParam) {
// CHECK-NEXT: CXXBoolLiteralExpr
// CHECK-NEXT: NullStmt
+#pragma acc parallel firstprivate(t, u)
+ while(true);
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+ // CHECK-NEXT: firstprivate clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'T' lvalue ParmVar{{.*}} 't' 'T'
+ // CHECK-NEXT: DeclRefExpr{{.*}}'U' lvalue ParmVar{{.*}} 'u' 'U'
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: NullStmt
+
#pragma acc parallel private(t) private(u)
while(true);
// CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
@@ -143,6 +191,18 @@ void TemplUses(T t, U u, T*PointerParam) {
// CHECK-NEXT: CXXBoolLiteralExpr
// CHECK-NEXT: NullStmt
+#pragma acc parallel private(t) firstprivate(NTTP, u)
+ while(true);
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+ // CHECK-NEXT: private clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'T' lvalue ParmVar{{.*}} 't' 'T'
+ // CHECK-NEXT: firstprivate clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'auto' lvalue NonTypeTemplateParm{{.*}} 'NTTP' 'auto &'
+ // CHECK-NEXT: DeclRefExpr{{.*}}'U' lvalue ParmVar{{.*}} 'u' 'U'
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: NullStmt
+
#pragma acc parallel private(u[0])
while(true);
// CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
@@ -206,6 +266,15 @@ void TemplUses(T t, U u, T*PointerParam) {
// CHECK-NEXT: CXXBoolLiteralExpr
// CHECK-NEXT: NullStmt
+// #pragma acc parallel firstprivate(t, u)
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+ // CHECK-NEXT: firstprivate clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue ParmVar{{.*}} 't' 'int'
+ // CHECK-NEXT: DeclRefExpr{{.*}}'int *' lvalue ParmVar{{.*}} 'u' 'int *'
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: NullStmt
+
// #pragma acc parallel private(t) private(u)
// CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
// CHECK-NEXT: private clause
@@ -229,6 +298,19 @@ void TemplUses(T t, U u, T*PointerParam) {
// CHECK-NEXT: CXXBoolLiteralExpr
// CHECK-NEXT: NullStmt
+// #pragma acc parallel private(t) firstprivate(NTTP, u)
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+ // CHECK-NEXT: private clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue ParmVar{{.*}} 't' 'int'
+ // CHECK-NEXT: firstprivate clause
+ // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}}'const unsigned int' lvalue
+ // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} referenced 'auto &' depth 0 index 0 NTTP
+ // CHECK-NEXT: DeclRefExpr{{.*}}'const unsigned int' lvalue Var{{.*}} 'CEVar' 'const unsigned int'
+ // CHECK-NEXT: DeclRefExpr{{.*}}'int *' lvalue ParmVar{{.*}} 'u' 'int *'
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: NullStmt
+
// #pragma acc parallel private(u[0])
// CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
// CHECK-NEXT: private clause
diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
index eb0ba09c5b9116..3f67fb93208954 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -2816,6 +2816,10 @@ void OpenACCClauseEnqueue::VisitNumGangsClause(const OpenACCNumGangsClause &C) {
void OpenACCClauseEnqueue::VisitPrivateClause(const OpenACCPrivateClause &C) {
VisitVarList(C);
}
+void OpenACCClauseEnqueue::VisitFirstPrivateClause(
+ const OpenACCFirstPrivateClause &C) {
+ VisitVarList(C);
+}
} // namespace
void EnqueueVisitor::EnqueueChildren(const OpenACCClause *C) {
More information about the cfe-commits
mailing list