Author: Erich Keane Date: 2024-10-15T06:12:19-07:00 New Revision: c8cbdc659c5f8044b5936ec40a4ab6a851ace2b3
URL: https://github.com/llvm/llvm-project/commit/c8cbdc659c5f8044b5936ec40a4ab6a851ace2b3 DIFF: https://github.com/llvm/llvm-project/commit/c8cbdc659c5f8044b5936ec40a4ab6a851ace2b3.diff LOG: [OpenACC] Implement 'loop' 'vector' clause (#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'. Added: clang/test/SemaOpenACC/loop-construct-vector-ast.cpp clang/test/SemaOpenACC/loop-construct-vector-clause.cpp Modified: clang/include/clang/AST/OpenACCClause.h clang/include/clang/Basic/DiagnosticSemaKinds.td 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/Sema/SemaOpenACC.cpp clang/lib/Sema/TreeTransform.h clang/lib/Serialization/ASTReader.cpp clang/lib/Serialization/ASTWriter.cpp clang/test/AST/ast-print-openacc-loop-construct.cpp clang/test/ParserOpenACC/parse-clauses.c clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c clang/test/SemaOpenACC/loop-construct-device_type-clause.c clang/tools/libclang/CIndex.cpp Removed: ################################################################################ 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@+1{{OpenACC clause 'vector' not yet implemented, clause ignored}} #pragma acc loop vector for(;;); // expected-error@+1{{expected expression}} @@ -908,8 +907,7 @@ void IntExprParsing() { // expected-error@+1{{expected expression}} #pragma acc loop vector(invalid:) for(;;); - // expected-error@+2{{invalid tag 'invalid' on 'vector' clause}} - // expected-warning@+1{{OpenACC clause 'vector' not yet implemented, clause ignored}} + // expected-error@+1{{invalid tag 'invalid' on 'vector' clause}} #pragma acc loop vector(invalid:5) for(;;); // expected-error@+1{{expected expression}} @@ -932,20 +930,15 @@ void IntExprParsing() { // expected-note@+1{{to match this '('}} #pragma acc loop vector(num:6,4) for(;;); - // expected-warning@+1{{OpenACC clause 'vector' not yet implemented, clause ignored}} #pragma acc loop vector(5) for(;;); - // expected-error@+2{{invalid tag 'num' on 'vector' clause}} - // expected-warning@+1{{OpenACC clause 'vector' not yet implemented, clause ignored}} + // expected-error@+1{{invalid tag 'num' on 'vector' clause}} #pragma acc loop vector(num:5) for(;;); - // expected-warning@+1{{OpenACC clause 'vector' not yet implemented, clause ignored}} #pragma acc loop vector(length:5) for(;;); - // expected-warning@+1{{OpenACC clause 'vector' not yet implemented, clause ignored}} #pragma acc loop vector(returns_int()) for(;;); - // expected-warning@+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@+1{{OpenACC clause 'vector' not yet implemented}} #pragma acc loop auto vector for(;;); // expected-warning@+1{{OpenACC clause 'nohost' not yet implemented}} @@ -181,7 +180,6 @@ void uses() { for(;;); #pragma acc loop worker auto for(;;); - // expected-warning@+1{{OpenACC clause 'vector' not yet implemented}} #pragma acc loop vector auto for(;;); // expected-warning@+1{{OpenACC clause 'nohost' not yet implemented}} @@ -318,7 +316,6 @@ void uses() { for(;;); #pragma acc loop independent worker for(;;); - // expected-warning@+1{{OpenACC clause 'vector' not yet implemented}} #pragma acc loop independent vector for(;;); // expected-warning@+1{{OpenACC clause 'nohost' not yet implemented}} @@ -454,7 +451,6 @@ void uses() { for(;;); #pragma acc loop worker independent for(;;); - // expected-warning@+1{{OpenACC clause 'vector' not yet implemented}} #pragma acc loop vector independent for(;;); // expected-warning@+1{{OpenACC clause 'nohost' not yet implemented}} @@ -591,9 +587,8 @@ void uses() { // expected-note@+1{{previous clause is here}} #pragma acc loop seq worker for(;;); - // expected-error@+3{{OpenACC clause 'vector' may not appear on the same construct as a 'seq' clause on a 'loop' construct}} - // expected-note@+2{{previous clause is here}} - // expected-warning@+1{{OpenACC clause 'vector' not yet implemented}} + // expected-error@+2{{OpenACC clause 'vector' may not appear on the same construct as a 'seq' clause on a 'loop' construct}} + // expected-note@+1{{previous clause is here}} #pragma acc loop seq vector for(;;); // expected-warning@+1{{OpenACC clause 'finalize' not yet implemented}} @@ -733,10 +728,8 @@ void uses() { // expected-note@+1{{previous clause is here}} #pragma acc loop worker seq for(;;); - // TODO OpenACC: when 'vector' is implemented and makes it to the AST, this should diagnose because of a conflict with 'seq'. - // TODOexpected-error@+3{{OpenACC clause 'vector' may not appear on the same construct as a 'seq' clause on a 'loop' construct}} - // TODOexpected-note@+2{{previous clause is here}} - // expected-warning@+1{{OpenACC clause 'vector' not yet implemented}} + // expected-error@+2{{OpenACC clause 'seq' may not appear on the same construct as a 'vector' clause on a 'loop' construct}} + // expected-note@+1{{previous clause is here}} #pragma acc loop vector seq for(;;); // expected-warning@+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@+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@+2{{OpenACC clause 'vector' requires expression of integer type ('NoConvert' invalid)}} +#pragma acc kernels +#pragma acc loop vector(length: NI) + for(;;); + + // expected-error@+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@+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@+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@+2{{loop with a 'vector' clause may not exist in the region of a 'vector' clause}} + // expected-note@-4{{previous clause is here}} +#pragma acc loop vector + for(;;); + for(;;); + } + +#pragma acc loop vector + for(;;) { + for(;;); + // expected-error@+4{{loop with a 'vector' clause may not exist in the region of a 'vector' clause}} + // expected-error@+3{{loop with a 'worker' clause may not exist in the region of a 'vector' clause}} + // expected-error@+2{{loop with a 'gang' clause may not exist in the region of a 'vector' clause}} + // expected-note@-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@+2{{OpenACC clause 'vector' requires expression of integer type ('NoConvert' invalid)}} +#pragma acc kernels +#pragma acc loop vector(length: NI) + for(;;); + + // expected-error@+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@+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@+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@+2{{loop with a 'vector' clause may not exist in the region of a 'vector' clause}} + // expected-note@-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@+4{{loop with a 'vector' clause may not exist in the region of a 'vector' clause}} + // expected-error@+3{{loop with a 'worker' clause may not exist in the region of a 'vector' clause}} + // expected-error@+2{{loop with a 'gang' clause may not exist in the region of a 'vector' clause}} + // expected-note@-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); _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
