Author: Erich Keane Date: 2024-10-01T06:40:21-07:00 New Revision: 97da34e0157d928e3cd0e9722b40ccf0d5769b5b
URL: https://github.com/llvm/llvm-project/commit/97da34e0157d928e3cd0e9722b40ccf0d5769b5b DIFF: https://github.com/llvm/llvm-project/commit/97da34e0157d928e3cd0e9722b40ccf0d5769b5b.diff LOG: [OpenACC] Add 'collapse' clause AST/basic Sema implementation (#109461) The 'collapse' clause on a 'loop' construct is used to specify how many nested loops are associated with the 'loop' construct. It takes an optional 'force' tag, and an integer constant expression as arguments. There are many other restrictions based on the contents of the loop/etc, but those are implemented in followup patches, for now, this patch just adds the AST node and does basic argument checking on the loop-count. Added: clang/test/SemaOpenACC/loop-construct-collapse-ast.cpp clang/test/SemaOpenACC/loop-construct-collapse-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/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-loop-construct.cpp clang/test/ParserOpenACC/parse-clauses.c clang/test/ParserOpenACC/parse-clauses.cpp clang/test/SemaOpenACC/compute-construct-device_type-clause.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 ea1ffbc7fd08b4..90f5b7fc9ab6f4 100644 --- a/clang/include/clang/AST/OpenACCClause.h +++ b/clang/include/clang/AST/OpenACCClause.h @@ -547,6 +547,32 @@ class OpenACCAsyncClause : public OpenACCClauseWithSingleIntExpr { SourceLocation EndLoc); }; +/// Represents a 'collapse' clause on a 'loop' construct. This clause takes an +/// integer constant expression 'N' that represents how deep to collapse the +/// construct. It also takes an optional 'force' tag that permits intervening +/// code in the loops. +class OpenACCCollapseClause : public OpenACCClauseWithSingleIntExpr { + bool HasForce = false; + + OpenACCCollapseClause(SourceLocation BeginLoc, SourceLocation LParenLoc, + bool HasForce, Expr *LoopCount, SourceLocation EndLoc); + +public: + const Expr *getLoopCount() const { return getIntExpr(); } + Expr *getLoopCount() { return getIntExpr(); } + + bool hasForce() const { return HasForce; } + + static bool classof(const OpenACCClause *C) { + return C->getClauseKind() == OpenACCClauseKind::Collapse; + } + + static OpenACCCollapseClause *Create(const ASTContext &C, + SourceLocation BeginLoc, + SourceLocation LParenLoc, bool HasForce, + Expr *LoopCount, SourceLocation EndLoc); +}; + /// Represents a clause with one or more 'var' objects, represented as an expr, /// as its arguments. Var-list is expected to be stored in trailing storage. /// For now, we're just storing the original expression in its entirety, unlike diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 0f591022e68541..64e6d0407b0ce3 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12606,6 +12606,9 @@ def note_acc_construct_here : Note<"'%0' construct is here">; def err_acc_loop_spec_conflict : Error<"OpenACC clause '%0' on '%1' construct conflicts with previous " "data dependence clause">; +def err_acc_collapse_loop_count + : Error<"OpenACC 'collapse' clause loop count must be a %select{constant " + "expression|positive integer value, evaluated to %1}0">; // AMDGCN builtins diagnostics def err_amdgcn_global_load_lds_size_invalid_value : Error<"invalid size value">; diff --git a/clang/include/clang/Basic/OpenACCClauses.def b/clang/include/clang/Basic/OpenACCClauses.def index 85f4859925f0bc..19cdfe7672133b 100644 --- a/clang/include/clang/Basic/OpenACCClauses.def +++ b/clang/include/clang/Basic/OpenACCClauses.def @@ -24,6 +24,7 @@ VISIT_CLAUSE(Auto) VISIT_CLAUSE(Async) VISIT_CLAUSE(Attach) +VISIT_CLAUSE(Collapse) VISIT_CLAUSE(Copy) CLAUSE_ALIAS(PCopy, Copy, true) CLAUSE_ALIAS(PresentOrCopy, Copy, true) diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h index 0ca76842e5f902..839fdb79cd0ac8 100644 --- a/clang/include/clang/Sema/SemaOpenACC.h +++ b/clang/include/clang/Sema/SemaOpenACC.h @@ -87,9 +87,14 @@ class SemaOpenACC : public SemaBase { SmallVector<Expr *> VarList; }; + struct CollapseDetails { + bool IsForce; + Expr *LoopCount; + }; + std::variant<std::monostate, DefaultDetails, ConditionDetails, IntExprDetails, VarListDetails, WaitDetails, DeviceTypeDetails, - ReductionDetails> + ReductionDetails, CollapseDetails> Details = std::monostate{}; public: @@ -246,6 +251,18 @@ class SemaOpenACC : public SemaBase { return std::get<VarListDetails>(Details).IsZero; } + bool isForce() const { + assert(ClauseKind == OpenACCClauseKind::Collapse && + "Only 'collapse' has a force tag"); + return std::get<CollapseDetails>(Details).IsForce; + } + + Expr *getLoopCount() const { + assert(ClauseKind == OpenACCClauseKind::Collapse && + "Only 'collapse' has a loop count"); + return std::get<CollapseDetails>(Details).LoopCount; + } + ArrayRef<DeviceTypeArgument> getDeviceTypeArchitectures() const { assert((ClauseKind == OpenACCClauseKind::DeviceType || ClauseKind == OpenACCClauseKind::DType) && @@ -384,6 +401,12 @@ class SemaOpenACC : public SemaBase { "Only 'device_type'/'dtype' has a device-type-arg list"); Details = DeviceTypeDetails{std::move(Archs)}; } + + void setCollapseDetails(bool IsForce, Expr *LoopCount) { + assert(ClauseKind == OpenACCClauseKind::Collapse && + "Only 'collapse' has collapse details"); + Details = CollapseDetails{IsForce, LoopCount}; + } }; SemaOpenACC(Sema &S); @@ -448,6 +471,8 @@ class SemaOpenACC : public SemaBase { Expr *LowerBound, SourceLocation ColonLocFirst, Expr *Length, SourceLocation RBLoc); + /// Checks the loop depth value for a collapse clause. + ExprResult CheckCollapseLoopCount(Expr *LoopCount); /// Helper type for the registration/assignment of constructs that need to /// 'know' about their parent constructs and hold a reference to them, such as diff --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp index 95089a9b79e267..d864ded33e8d1f 100644 --- a/clang/lib/AST/OpenACCClause.cpp +++ b/clang/lib/AST/OpenACCClause.cpp @@ -43,7 +43,7 @@ bool OpenACCClauseWithCondition::classof(const OpenACCClause *C) { bool OpenACCClauseWithSingleIntExpr::classof(const OpenACCClause *C) { return OpenACCNumWorkersClause::classof(C) || OpenACCVectorLengthClause::classof(C) || - OpenACCAsyncClause::classof(C); + OpenACCCollapseClause::classof(C) || OpenACCAsyncClause::classof(C); } OpenACCDefaultClause *OpenACCDefaultClause::Create(const ASTContext &C, OpenACCDefaultClauseKind K, @@ -134,6 +134,30 @@ OpenACCNumWorkersClause::Create(const ASTContext &C, SourceLocation BeginLoc, OpenACCNumWorkersClause(BeginLoc, LParenLoc, IntExpr, EndLoc); } +OpenACCCollapseClause::OpenACCCollapseClause(SourceLocation BeginLoc, + SourceLocation LParenLoc, + bool HasForce, Expr *LoopCount, + SourceLocation EndLoc) + : OpenACCClauseWithSingleIntExpr(OpenACCClauseKind::Collapse, BeginLoc, + LParenLoc, LoopCount, EndLoc), + HasForce(HasForce) { + assert(LoopCount && "LoopCount required"); +} + +OpenACCCollapseClause * +OpenACCCollapseClause::Create(const ASTContext &C, SourceLocation BeginLoc, + SourceLocation LParenLoc, bool HasForce, + Expr *LoopCount, SourceLocation EndLoc) { + assert( + LoopCount && + (LoopCount->isInstantiationDependent() || isa<ConstantExpr>(LoopCount)) && + "Loop count not constant expression"); + void *Mem = + C.Allocate(sizeof(OpenACCCollapseClause), alignof(OpenACCCollapseClause)); + return new (Mem) + OpenACCCollapseClause(BeginLoc, LParenLoc, HasForce, LoopCount, EndLoc); +} + OpenACCVectorLengthClause::OpenACCVectorLengthClause(SourceLocation BeginLoc, SourceLocation LParenLoc, Expr *IntExpr, @@ -550,3 +574,11 @@ void OpenACCClausePrinter::VisitIndependentClause( void OpenACCClausePrinter::VisitSeqClause(const OpenACCSeqClause &C) { OS << "seq"; } + +void OpenACCClausePrinter::VisitCollapseClause(const OpenACCCollapseClause &C) { + OS << "collapse("; + if (C.hasForce()) + OS << "force:"; + printExpr(C.getLoopCount()); + OS << ")"; +} diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp index ad4281986f668e..c3812844ab8a31 100644 --- a/clang/lib/AST/StmtProfile.cpp +++ b/clang/lib/AST/StmtProfile.cpp @@ -2558,6 +2558,12 @@ void OpenACCClauseProfiler::VisitNumWorkersClause( Profiler.VisitStmt(Clause.getIntExpr()); } +void OpenACCClauseProfiler::VisitCollapseClause( + const OpenACCCollapseClause &Clause) { + assert(Clause.getLoopCount() && "collapse clause requires a valid int expr"); + Profiler.VisitStmt(Clause.getLoopCount()); +} + void OpenACCClauseProfiler::VisitPrivateClause( const OpenACCPrivateClause &Clause) { for (auto *E : Clause.getVarList()) diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp index 3c51c746471829..8a74159c7c93e5 100644 --- a/clang/lib/AST/TextNodeDumper.cpp +++ b/clang/lib/AST/TextNodeDumper.cpp @@ -419,6 +419,12 @@ void TextNodeDumper::Visit(const OpenACCClause *C) { // but print 'clause' here so it is clear what is happening from the dump. OS << " clause"; break; + case OpenACCClauseKind::Collapse: + OS << " clause"; + if (cast<OpenACCCollapseClause>(C)->hasForce()) + OS << ": force"; + break; + case OpenACCClauseKind::CopyIn: case OpenACCClauseKind::PCopyIn: case OpenACCClauseKind::PresentOrCopyIn: diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp index 0261e8ea3c9b76..e66abd6873794e 100644 --- a/clang/lib/Parse/ParseOpenACC.cpp +++ b/clang/lib/Parse/ParseOpenACC.cpp @@ -976,14 +976,25 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams( /*IsReadOnly=*/false, /*IsZero=*/false); break; case OpenACCClauseKind::Collapse: { - tryParseAndConsumeSpecialTokenKind(*this, OpenACCSpecialTokenKind::Force, - ClauseKind); - ExprResult NumLoops = + bool HasForce = tryParseAndConsumeSpecialTokenKind( + *this, OpenACCSpecialTokenKind::Force, ClauseKind); + ExprResult LoopCount = getActions().CorrectDelayedTyposInExpr(ParseConstantExpression()); - if (NumLoops.isInvalid()) { + if (LoopCount.isInvalid()) { Parens.skipToEnd(); return OpenACCCanContinue(); } + + LoopCount = getActions().OpenACC().ActOnIntExpr( + OpenACCDirectiveKind::Invalid, ClauseKind, + LoopCount.get()->getBeginLoc(), LoopCount.get()); + + if (LoopCount.isInvalid()) { + Parens.skipToEnd(); + return OpenACCCanContinue(); + } + + ParsedClause.setCollapseDetails(HasForce, LoopCount.get()); break; } case OpenACCClauseKind::Bind: { diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp index ecbcc19413dc61..89142b837e60a9 100644 --- a/clang/lib/Sema/SemaOpenACC.cpp +++ b/clang/lib/Sema/SemaOpenACC.cpp @@ -343,6 +343,18 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind, return false; } + case OpenACCClauseKind::Collapse: { + switch (DirectiveKind) { + case OpenACCDirectiveKind::Loop: + case OpenACCDirectiveKind::ParallelLoop: + case OpenACCDirectiveKind::SerialLoop: + case OpenACCDirectiveKind::KernelsLoop: + return true; + default: + return false; + } + } + default: // Do nothing so we can go to the 'unimplemented' diagnostic instead. return true; @@ -1037,6 +1049,26 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitReductionClause( ValidVars, Clause.getEndLoc()); } +OpenACCClause *SemaOpenACCClauseVisitor::VisitCollapseClause( + SemaOpenACC::OpenACCParsedClause &Clause) { + // Duplicates here are not really sensible. We could possible permit + // multiples if they all had the same value, but there isn't really a good + // reason to do so. Also, this simplifies the suppression of duplicates, in + // that we know if we 'find' one after instantiation, that it is the same + // clause, which simplifies instantiation/checking/etc. + if (checkAlreadyHasClauseOfKind(SemaRef, ExistingClauses, Clause)) + return nullptr; + + ExprResult LoopCount = SemaRef.CheckCollapseLoopCount(Clause.getLoopCount()); + + if (!LoopCount.isUsable()) + return nullptr; + + return OpenACCCollapseClause::Create(Ctx, Clause.getBeginLoc(), + Clause.getLParenLoc(), Clause.isForce(), + LoopCount.get(), Clause.getEndLoc()); +} + } // namespace SemaOpenACC::SemaOpenACC(Sema &S) : SemaBase(S) {} @@ -1273,6 +1305,9 @@ ExprResult SemaOpenACC::ActOnIntExpr(OpenACCDirectiveKind DK, } } IntExprDiagnoser(DK, CK, IntExpr); + if (!IntExpr) + return ExprError(); + ExprResult IntExprResult = SemaRef.PerformContextualImplicitConversion( Loc, IntExpr, IntExprDiagnoser); if (IntExprResult.isInvalid()) @@ -1583,6 +1618,34 @@ ExprResult SemaOpenACC::ActOnArraySectionExpr(Expr *Base, SourceLocation LBLoc, OK_Ordinary, ColonLoc, RBLoc); } +ExprResult SemaOpenACC::CheckCollapseLoopCount(Expr *LoopCount) { + if (!LoopCount) + return ExprError(); + + assert((LoopCount->isInstantiationDependent() || + LoopCount->getType()->isIntegerType()) && + "Loop argument non integer?"); + + // If this is dependent, there really isn't anything we can check. + if (LoopCount->isInstantiationDependent()) + return ExprResult{LoopCount}; + + std::optional<llvm::APSInt> ICE = + LoopCount->getIntegerConstantExpr(getASTContext()); + + // OpenACC 3.3: 2.9.1 + // The argument to the collapse clause must be a constant positive integer + // expression. + if (!ICE || *ICE <= 0) { + Diag(LoopCount->getBeginLoc(), diag::err_acc_collapse_loop_count) + << ICE.has_value() << ICE.value_or(llvm::APSInt{}).getExtValue(); + return ExprError(); + } + + return ExprResult{ + ConstantExpr::Create(getASTContext(), LoopCount, APValue{*ICE})}; +} + bool SemaOpenACC::ActOnStartStmtDirective(OpenACCDirectiveKind K, SourceLocation StartLoc) { SemaRef.DiscardCleanupsInEvaluationContext(); diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h index 91cb980ee26b26..6fdb18d51acef9 100644 --- a/clang/lib/Sema/TreeTransform.h +++ b/clang/lib/Sema/TreeTransform.h @@ -11822,6 +11822,31 @@ void OpenACCClauseTransform<Derived>::VisitReductionClause( ParsedClause.getLParenLoc(), C.getReductionOp(), ValidVars, ParsedClause.getEndLoc()); } + +template <typename Derived> +void OpenACCClauseTransform<Derived>::VisitCollapseClause( + const OpenACCCollapseClause &C) { + Expr *LoopCount = const_cast<Expr *>(C.getLoopCount()); + assert(LoopCount && "collapse clause constructed with invalid loop count"); + + ExprResult NewLoopCount = Self.TransformExpr(LoopCount); + + NewLoopCount = Self.getSema().OpenACC().ActOnIntExpr( + OpenACCDirectiveKind::Invalid, ParsedClause.getClauseKind(), + NewLoopCount.get()->getBeginLoc(), NewLoopCount.get()); + + NewLoopCount = + Self.getSema().OpenACC().CheckCollapseLoopCount(NewLoopCount.get()); + + if (!NewLoopCount.isUsable()) + return; + + ParsedClause.setCollapseDetails(C.hasForce(), NewLoopCount.get()); + NewClause = OpenACCCollapseClause::Create( + Self.getSema().getASTContext(), ParsedClause.getBeginLoc(), + ParsedClause.getLParenLoc(), ParsedClause.isForce(), + ParsedClause.getLoopCount(), ParsedClause.getEndLoc()); +} } // namespace template <typename Derived> OpenACCClause *TreeTransform<Derived>::TransformOpenACCClause( diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp index d0dff9a354c108..0a4251c0e52404 100644 --- a/clang/lib/Serialization/ASTReader.cpp +++ b/clang/lib/Serialization/ASTReader.cpp @@ -12283,6 +12283,13 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() { return OpenACCIndependentClause::Create(getContext(), BeginLoc, EndLoc); case OpenACCClauseKind::Auto: return OpenACCAutoClause::Create(getContext(), BeginLoc, EndLoc); + case OpenACCClauseKind::Collapse: { + SourceLocation LParenLoc = readSourceLocation(); + bool HasForce = readBool(); + Expr *LoopCount = readSubExpr(); + return OpenACCCollapseClause::Create(getContext(), BeginLoc, LParenLoc, + HasForce, LoopCount, EndLoc); + } case OpenACCClauseKind::Finalize: case OpenACCClauseKind::IfPresent: @@ -12296,7 +12303,6 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() { case OpenACCClauseKind::DeviceResident: case OpenACCClauseKind::Host: case OpenACCClauseKind::Link: - case OpenACCClauseKind::Collapse: case OpenACCClauseKind::Bind: case OpenACCClauseKind::DeviceNum: case OpenACCClauseKind::DefaultAsync: diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp index 7a40c5c65d39d9..aa9764e25c3233 100644 --- a/clang/lib/Serialization/ASTWriter.cpp +++ b/clang/lib/Serialization/ASTWriter.cpp @@ -8148,6 +8148,13 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) { // Nothing to do here, there is no additional information beyond the // begin/end loc and clause kind. return; + case OpenACCClauseKind::Collapse: { + const auto *CC = cast<OpenACCCollapseClause>(C); + writeSourceLocation(CC->getLParenLoc()); + writeBool(CC->hasForce()); + AddStmt(const_cast<Expr *>(CC->getLoopCount())); + return; + } case OpenACCClauseKind::Finalize: case OpenACCClauseKind::IfPresent: @@ -8161,7 +8168,6 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) { case OpenACCClauseKind::DeviceResident: case OpenACCClauseKind::Host: case OpenACCClauseKind::Link: - case OpenACCClauseKind::Collapse: case OpenACCClauseKind::Bind: case OpenACCClauseKind::DeviceNum: case OpenACCClauseKind::DefaultAsync: diff --git a/clang/test/AST/ast-print-openacc-loop-construct.cpp b/clang/test/AST/ast-print-openacc-loop-construct.cpp index cde302a66f3af7..ae1f7964f019eb 100644 --- a/clang/test/AST/ast-print-openacc-loop-construct.cpp +++ b/clang/test/AST/ast-print-openacc-loop-construct.cpp @@ -57,4 +57,29 @@ void foo() { // CHECK-NEXT: ; #pragma acc loop private(i, array[1], array, array[1:2]) for(;;); + +// CHECK: #pragma acc loop collapse(1) +// CHECK-NEXT: for (;;) +// CHECK-NEXT: ; +#pragma acc loop collapse(1) + for(;;); +// CHECK: #pragma acc loop collapse(force:1) +// CHECK-NEXT: for (;;) +// CHECK-NEXT: ; +#pragma acc loop collapse(force:1) + for(;;); +// CHECK: #pragma acc loop collapse(2) +// CHECK-NEXT: for (;;) +// CHECK-NEXT: for (;;) +// CHECK-NEXT: ; +#pragma acc loop collapse(2) + for(;;) + for(;;); +// CHECK: #pragma acc loop collapse(force:2) +// CHECK-NEXT: for (;;) +// CHECK-NEXT: for (;;) +// CHECK-NEXT: ; +#pragma acc loop collapse(force:2) + for(;;) + for(;;); } diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c index 15c4554a31922a..6c9ce4ad5e1969 100644 --- a/clang/test/ParserOpenACC/parse-clauses.c +++ b/clang/test/ParserOpenACC/parse-clauses.c @@ -105,17 +105,14 @@ void func() { #pragma acc loop collapse(force:) for(;;){} - // expected-error@+2{{invalid tag 'unknown' on 'collapse' clause}} - // expected-warning@+1{{OpenACC clause 'collapse' not yet implemented, clause ignored}} -#pragma acc loop collapse(unknown:5) + // expected-error@+1{{invalid tag 'unknown' on 'collapse' clause}} +#pragma acc loop collapse(unknown:1) for(;;){} - // expected-warning@+1{{OpenACC clause 'collapse' not yet implemented, clause ignored}} -#pragma acc loop collapse(force:5) +#pragma acc loop collapse(force:1) for(;;){} - // expected-warning@+1{{OpenACC clause 'collapse' not yet implemented, clause ignored}} -#pragma acc loop collapse(5) +#pragma acc loop collapse(1) for(;;){} // expected-error@+2{{expected ')'}} diff --git a/clang/test/ParserOpenACC/parse-clauses.cpp b/clang/test/ParserOpenACC/parse-clauses.cpp index b7e252e892beab..9613530db77ddc 100644 --- a/clang/test/ParserOpenACC/parse-clauses.cpp +++ b/clang/test/ParserOpenACC/parse-clauses.cpp @@ -2,13 +2,23 @@ template<unsigned I, typename T> void templ() { - // expected-warning@+1{{OpenACC clause 'collapse' not yet implemented, clause ignored}} #pragma acc loop collapse(I) - for(;;){} + for(;;) + for(;;) + for(;;) + for(;;) + for(;;) + for(;;) + for(;;); - // expected-warning@+1{{OpenACC clause 'collapse' not yet implemented, clause ignored}} #pragma acc loop collapse(T::value) - for(;;){} + for(;;) + for(;;) + for(;;) + for(;;) + for(;;) + for(;;) + for(;;); #pragma acc parallel vector_length(T::value) for(;;){} diff --git a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c index b300abe577801c..26f0315fb86f1a 100644 --- a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c +++ b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c @@ -188,8 +188,7 @@ void uses() { // expected-note@+1{{previous clause is here}} #pragma acc serial device_type(*) reduction(+:Var) while(1); - // expected-error@+2{{OpenACC clause 'collapse' may not follow a 'device_type' clause in a compute construct}} - // expected-note@+1{{previous clause is here}} + // expected-error@+1{{OpenACC 'collapse' clause is not valid on 'kernels' directive}} #pragma acc kernels device_type(*) collapse(1) while(1); // expected-error@+2{{OpenACC clause 'bind' may not follow a 'device_type' clause in a compute construct}} diff --git a/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c b/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c index ac61976ff620d8..3212c19d089fc9 100644 --- a/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c +++ b/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c @@ -138,7 +138,6 @@ void uses() { // expected-warning@+1{{OpenACC clause 'reduction' not yet implemented}} #pragma acc loop auto reduction(+:Var) for(;;); - // expected-warning@+1{{OpenACC clause 'collapse' not yet implemented}} #pragma acc loop auto collapse(1) for(;;); // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}} @@ -277,7 +276,6 @@ void uses() { // expected-warning@+1{{OpenACC clause 'reduction' not yet implemented}} #pragma acc loop reduction(+:Var) auto for(;;); - // expected-warning@+1{{OpenACC clause 'collapse' not yet implemented}} #pragma acc loop collapse(1) auto for(;;); // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}} @@ -417,7 +415,6 @@ void uses() { // expected-warning@+1{{OpenACC clause 'reduction' not yet implemented}} #pragma acc loop independent reduction(+:Var) for(;;); - // expected-warning@+1{{OpenACC clause 'collapse' not yet implemented}} #pragma acc loop independent collapse(1) for(;;); // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}} @@ -556,7 +553,6 @@ void uses() { // expected-warning@+1{{OpenACC clause 'reduction' not yet implemented}} #pragma acc loop reduction(+:Var) independent for(;;); - // expected-warning@+1{{OpenACC clause 'collapse' not yet implemented}} #pragma acc loop collapse(1) independent for(;;); // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}} @@ -705,7 +701,6 @@ void uses() { // expected-warning@+1{{OpenACC clause 'reduction' not yet implemented}} #pragma acc loop seq reduction(+:Var) for(;;); - // expected-warning@+1{{OpenACC clause 'collapse' not yet implemented}} #pragma acc loop seq collapse(1) for(;;); // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}} @@ -853,7 +848,6 @@ void uses() { // expected-warning@+1{{OpenACC clause 'reduction' not yet implemented}} #pragma acc loop reduction(+:Var) seq for(;;); - // expected-warning@+1{{OpenACC clause 'collapse' not yet implemented}} #pragma acc loop collapse(1) seq for(;;); // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}} diff --git a/clang/test/SemaOpenACC/loop-construct-collapse-ast.cpp b/clang/test/SemaOpenACC/loop-construct-collapse-ast.cpp new file mode 100644 index 00000000000000..3bdcfbf95b96c3 --- /dev/null +++ b/clang/test/SemaOpenACC/loop-construct-collapse-ast.cpp @@ -0,0 +1,158 @@ +// 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 + +struct S { + constexpr S(){}; + constexpr operator auto() {return 1;} +}; + +void NormalUses() { + // CHECK: FunctionDecl{{.*}}NormalUses + // CHECK-NEXT: CompoundStmt + +#pragma acc loop collapse(1) + for(;;); + // CHECK-NEXT: OpenACCLoopConstruct + // CHECK-NEXT: collapse clause + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1 + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: NullStmt + +#pragma acc loop collapse(force:S{}) + for(;;); + // CHECK-NEXT: OpenACCLoopConstruct + // CHECK-NEXT: collapse clause + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <UserDefinedConversion> + // CHECK-NEXT: CXXMemberCallExpr{{.*}} 'int' + // CHECK-NEXT: MemberExpr{{.*}} .operator auto + // CHECK-NEXT: MaterializeTemporaryExpr{{.*}}'S' lvalue + // CHECK-NEXT: CXXTemporaryObjectExpr{{.*}}'S' 'void ()' list + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: NullStmt +} + +template<typename T, unsigned Value> +void TemplUses() { + // CHECK: FunctionTemplateDecl{{.*}}TemplUses + // CHECK-NEXT: TemplateTypeParmDecl {{.*}} referenced typename depth 0 index 0 T + // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} referenced 'unsigned int' depth 0 index 1 Value + // CHECK-NEXT: FunctionDecl{{.*}} TemplUses 'void ()' + // CHECK-NEXT: CompoundStmt + +#pragma acc loop collapse(Value) + for(;;) + for(;;); + // CHECK-NEXT: OpenACCLoopConstruct + // CHECK-NEXT: collapse clause + // CHECK-NEXT: DeclRefExpr{{.*}} 'unsigned int' NonTypeTemplateParm{{.*}} 'Value' + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: NullStmt + +#pragma acc loop collapse(force:T{} + S{}) + for(;;) + for(;;); + // CHECK-NEXT: OpenACCLoopConstruct + // CHECK-NEXT: collapse clause + // CHECK-NEXT: BinaryOperator {{.*}}'+' + // CHECK-NEXT: CXXUnresolvedConstructExpr{{.*}}'T' 'T' list + // CHECK-NEXT: InitListExpr + // CHECK-NEXT: CXXTemporaryObjectExpr{{.*}}'S' 'void ()' list + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: NullStmt + + // Instantiation: + // CHECK-NEXT: FunctionDecl{{.*}} used TemplUses 'void ()' implicit_instantiation + // CHECK-NEXT: TemplateArgument type 'S' + // CHECK-NEXT: RecordType{{.*}} 'S' + // CHECK-NEXT: CXXRecord{{.*}} 'S' + // CHECK-NEXT: TemplateArgument integral '2U' + // CHECK-NEXT: CompoundStmt + + // CHECK-NEXT: OpenACCLoopConstruct + // CHECK-NEXT: collapse clause + // CHECK-NEXT: ConstantExpr{{.*}}'unsigned int' + // CHECK-NEXT: value: Int 2 + // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}} 'unsigned int' + // CHECK-NEXT: NonTypeTemplateParmDecl + // CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned int' 2 + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: NullStmt + + // CHECK-NEXT: OpenACCLoopConstruct + // CHECK-NEXT: collapse clause + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 2 + // CHECK-NEXT: BinaryOperator {{.*}}'+' + // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <UserDefinedConversion> + // CHECK-NEXT: CXXMemberCallExpr{{.*}} 'int' + // CHECK-NEXT: MemberExpr{{.*}} .operator auto + // CHECK-NEXT: MaterializeTemporaryExpr{{.*}}'S' lvalue + // CHECK-NEXT: CXXTemporaryObjectExpr{{.*}}'S' 'void ()' list + // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <UserDefinedConversion> + // CHECK-NEXT: CXXMemberCallExpr{{.*}} 'int' + // CHECK-NEXT: MemberExpr{{.*}} .operator auto + // CHECK-NEXT: MaterializeTemporaryExpr{{.*}}'S' lvalue + // CHECK-NEXT: CXXTemporaryObjectExpr{{.*}}'S' 'void ()' list + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: <<<NULL>>> + // CHECK-NEXT: NullStmt + +} + +void Inst() { + TemplUses<S, 2>(); +} + +#endif // PCH_HELPER diff --git a/clang/test/SemaOpenACC/loop-construct-collapse-clause.cpp b/clang/test/SemaOpenACC/loop-construct-collapse-clause.cpp new file mode 100644 index 00000000000000..9c1e577773e8f8 --- /dev/null +++ b/clang/test/SemaOpenACC/loop-construct-collapse-clause.cpp @@ -0,0 +1,117 @@ +// RUN: %clang_cc1 %s -fopenacc -verify + + +void only_for_loops() { + // expected-error@+3{{OpenACC 'loop' construct can only be applied to a 'for' loop}} + // expected-note@+1{{'loop' construct is here}} +#pragma acc loop collapse(1) + while(true); + + // expected-error@+3{{OpenACC 'loop' construct can only be applied to a 'for' loop}} + // expected-note@+1{{'loop' construct is here}} +#pragma acc loop collapse(1) + do{}while(true); + +} + +void only_one_on_loop() { + // expected-error@+2{{OpenACC 'collapse' clause cannot appear more than once on a 'loop' directive}} + // expected-note@+1{{previous clause is here}} +#pragma acc loop collapse(1) collapse(1) + for(;;); +} + +constexpr int three() { return 3; } +constexpr int one() { return 1; } +constexpr int neg() { return -1; } +constexpr int zero() { return 0; } + +struct NotConstexpr { + constexpr NotConstexpr(){}; + + operator int(){ return 1; } +}; +struct ConvertsNegative { + constexpr ConvertsNegative(){}; + + constexpr operator int(){ return -1; } +}; +struct ConvertsOne{ + constexpr ConvertsOne(){}; + + constexpr operator int(){ return 1; } +}; + +struct ConvertsThree{ + constexpr ConvertsThree(){}; + + constexpr operator int(){ return 3; } +}; + +template <typename T, int Val> +void negative_constexpr_templ() { + // expected-error@+3 2{{OpenACC 'collapse' clause loop count must be a positive integer value, evaluated to 0}} + // expected-note@#NCETN1{{in instantiation of function template specialization 'negative_constexpr_templ<int, -1>'}} + // expected-note@#NCET1{{in instantiation of function template specialization 'negative_constexpr_templ<int, 1>'}} +#pragma acc loop collapse(T{}) + for(;;) + for(;;); + + // expected-error@+1{{OpenACC 'collapse' clause loop count must be a positive integer value, evaluated to -1}} +#pragma acc loop collapse(Val) + for(;;) + for(;;); +} + +void negative_constexpr(int i) { +#pragma acc loop collapse(2) + for(;;) + for(;;); + +#pragma acc loop collapse(1) + for(;;) + for(;;); + + // expected-error@+1{{OpenACC 'collapse' clause loop count must be a positive integer value, evaluated to 0}} +#pragma acc loop collapse(0) + for(;;) + for(;;); + + // expected-error@+1{{OpenACC 'collapse' clause loop count must be a positive integer value, evaluated to -1}} +#pragma acc loop collapse(-1) + for(;;) + for(;;); + +#pragma acc loop collapse(one()) + for(;;) + for(;;); + + // expected-error@+1{{OpenACC 'collapse' clause loop count must be a positive integer value, evaluated to 0}} +#pragma acc loop collapse(zero()) + for(;;) + for(;;); + + // expected-error@+1{{OpenACC 'collapse' clause loop count must be a positive integer value, evaluated to -1}} +#pragma acc loop collapse(neg()) + for(;;) + for(;;); + + // expected-error@+1{{OpenACC 'collapse' clause loop count must be a constant expression}} +#pragma acc loop collapse(NotConstexpr{}) + for(;;) + for(;;); + + // expected-error@+1{{OpenACC 'collapse' clause loop count must be a positive integer value, evaluated to -1}} +#pragma acc loop collapse(ConvertsNegative{}) + for(;;) + for(;;); + +#pragma acc loop collapse(ConvertsOne{}) + for(;;) + for(;;); + + negative_constexpr_templ<int, -1>(); // #NCETN1 + + negative_constexpr_templ<int, 1>(); // #NCET1 +} + diff --git a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c index 520ba45aaebf4f..47c9239f4f0e96 100644 --- a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c +++ b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c @@ -162,7 +162,6 @@ void uses() { // expected-note@+1{{previous clause is here}} #pragma acc loop device_type(*) reduction(+:Var) for(;;); - // expected-warning@+1{{OpenACC clause 'collapse' not yet implemented, clause ignored}} #pragma acc loop device_type(*) collapse(1) for(;;); // expected-error@+2{{OpenACC clause 'bind' may not follow a 'device_type' clause in a 'loop' construct}} diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp index 9a2be8e3aabb62..d188f794bad20e 100644 --- a/clang/tools/libclang/CIndex.cpp +++ b/clang/tools/libclang/CIndex.cpp @@ -2893,6 +2893,9 @@ void OpenACCClauseEnqueue::VisitAutoClause(const OpenACCAutoClause &C) {} void OpenACCClauseEnqueue::VisitIndependentClause( const OpenACCIndependentClause &C) {} void OpenACCClauseEnqueue::VisitSeqClause(const OpenACCSeqClause &C) {} +void OpenACCClauseEnqueue::VisitCollapseClause(const OpenACCCollapseClause &C) { + Visitor.AddStmt(C.getLoopCount()); +} } // 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
