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@+3{{expected ','}} - // expected-warning@+2{{OpenACC clause 'firstprivate' not yet implemented, clause ignored}} + // expected-error@+2{{expected ','}} // expected-warning@+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@+2{{OpenACC clause 'firstprivate' not yet implemented, clause ignored}} // expected-warning@+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@+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@+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@+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@+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@+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@+2{{OpenACC sub-array specified range [2:5] would be out of the range of the subscripted array size of 5}} + // expected-error@+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@+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@+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@+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@+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@+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@+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@+2{{OpenACC sub-array specified range [2:5] would be out of the range of the subscripted array size of 5}} + // expected-error@+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@+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@+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@+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@+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@+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@+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) { _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
