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 ba813af960af6f..fb60477bc0c2e6 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 VarList; }; + struct CollapseDetails { + bool IsForce; + Expr *LoopCount; + }; + std::variant + ReductionDetails, CollapseDetails> Details = std::monostate{}; public: @@ -246,6 +251,18 @@ class SemaOpenACC : public SemaBase { return std::get(Details).IsZero; } + bool isForce() const { + assert(ClauseKind == OpenACCClauseKind::Collapse && + "Only 'collapse' has a force tag"); + return std::get(Details).IsForce; + } + + Expr *getLoopCount() const { + assert(ClauseKind == OpenACCClauseKind::Collapse && + "Only 'collapse' has a loop count"); + return std::get(Details).LoopCount; + } + ArrayRef 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(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(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 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 95ded5e59a9fa7..3dba8e53669e8d 100644 --- a/clang/lib/Sema/TreeTransform.h +++ b/clang/lib/Sema/TreeTransform.h @@ -11821,6 +11821,31 @@ void OpenACCClauseTransform::VisitReductionClause( ParsedClause.getLParenLoc(), C.getReductionOp(), ValidVars, ParsedClause.getEndLoc()); } + +template +void OpenACCClauseTransform::VisitCollapseClause( + const OpenACCCollapseClause &C) { + Expr *LoopCount = const_cast(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 OpenACCClause *TreeTransform::TransformOpenACCClause( diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp index 7efcc81e194d95..7a0a6d81b21a58 100644 --- a/clang/lib/Serialization/ASTReader.cpp +++ b/clang/lib/Serialization/ASTReader.cpp @@ -12269,6 +12269,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: @@ -12282,7 +12289,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 008bf571f847dc..2da0c94c9e4cbc 100644 --- a/clang/lib/Serialization/ASTWriter.cpp +++ b/clang/lib/Serialization/ASTWriter.cpp @@ -8126,6 +8126,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(C); + writeSourceLocation(CC->getLParenLoc()); + writeBool(CC->hasForce()); + AddStmt(const_cast(CC->getLoopCount())); + return; + } case OpenACCClauseKind::Finalize: case OpenACCClauseKind::IfPresent: @@ -8139,7 +8146,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 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: <<>> + // CHECK-NEXT: <<>> + // CHECK-NEXT: <<>> + // CHECK-NEXT: <<>> + // 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' + // 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: <<>> + // CHECK-NEXT: <<>> + // CHECK-NEXT: <<>> + // CHECK-NEXT: <<>> + // CHECK-NEXT: NullStmt +} + +template +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: <<>> + // CHECK-NEXT: <<>> + // CHECK-NEXT: <<>> + // CHECK-NEXT: <<>> + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<>> + // CHECK-NEXT: <<>> + // CHECK-NEXT: <<>> + // CHECK-NEXT: <<>> + // 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: <<>> + // CHECK-NEXT: <<>> + // CHECK-NEXT: <<>> + // CHECK-NEXT: <<>> + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<>> + // CHECK-NEXT: <<>> + // CHECK-NEXT: <<>> + // CHECK-NEXT: <<>> + // 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: <<>> + // CHECK-NEXT: <<>> + // CHECK-NEXT: <<>> + // CHECK-NEXT: <<>> + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<>> + // CHECK-NEXT: <<>> + // CHECK-NEXT: <<>> + // CHECK-NEXT: <<>> + // 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' + // CHECK-NEXT: CXXMemberCallExpr{{.*}} 'int' + // CHECK-NEXT: MemberExpr{{.*}} .operator auto + // CHECK-NEXT: MaterializeTemporaryExpr{{.*}}'S' lvalue + // CHECK-NEXT: CXXTemporaryObjectExpr{{.*}}'S' 'void ()' list + // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' + // 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: <<>> + // CHECK-NEXT: <<>> + // CHECK-NEXT: <<>> + // CHECK-NEXT: <<>> + // CHECK-NEXT: ForStmt + // CHECK-NEXT: <<>> + // CHECK-NEXT: <<>> + // CHECK-NEXT: <<>> + // CHECK-NEXT: <<>> + // CHECK-NEXT: NullStmt + +} + +void Inst() { + TemplUses(); +} + +#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 +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'}} + // expected-note@#NCET1{{in instantiation of function template specialization 'negative_constexpr_templ'}} +#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(); // #NCETN1 + + negative_constexpr_templ(); // #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 e821c5e4c588b6..83a015b8f3ca71 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) {