Author: erichkeane Date: 2024-12-09T14:06:44-08:00 New Revision: 7d89ebfd5f93577e7b1f12d1d21ee3e87eacde07
URL: https://github.com/llvm/llvm-project/commit/7d89ebfd5f93577e7b1f12d1d21ee3e87eacde07 DIFF: https://github.com/llvm/llvm-project/commit/7d89ebfd5f93577e7b1f12d1d21ee3e87eacde07.diff LOG: [OpenACC] Implement 'reduction' for combined constructs. Once again, this is a clause on a combined construct that does almost exactly what the loop/compute construct version does, only with some sl ightly different evaluation rules/sema rules as it doesn't have to consider the parent, just the 'combined' construct. The two sets of rules for reduction on loop and compute are fine together, so this ensures they are all enforced for this too. The 'gangs' 'num_gangs' 'reduction' diagnostic (Dim>1) had to be applied to num_gangs as well, as it previously wasn't permissible to get in this situation, but we now can. Added: clang/test/SemaOpenACC/combined-construct-reduction-ast.cpp clang/test/SemaOpenACC/combined-construct-reduction-clause.cpp Modified: clang/include/clang/Basic/DiagnosticSemaKinds.td clang/include/clang/Sema/SemaOpenACC.h clang/lib/Sema/SemaOpenACC.cpp clang/lib/Sema/TreeTransform.h clang/test/AST/ast-print-openacc-combined-construct.cpp clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c clang/test/SemaOpenACC/compute-construct-reduction-clause.c clang/test/SemaOpenACC/compute-construct-reduction-clause.cpp Removed: ################################################################################ diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index f185a98720f5b9..0a245e2077f68f 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12716,9 +12716,9 @@ def err_acc_clause_cannot_combine : Error<"OpenACC clause '%0' may not appear on the same construct as a " "'%1' clause on a '%2' construct">; def err_acc_reduction_num_gangs_conflict - : Error< - "OpenACC 'reduction' clause may not appear on a 'parallel' construct " - "with a 'num_gangs' clause with more than 1 argument, have %0">; + : Error<"OpenACC '%1' clause %select{|with more than 1 argument }0may not " + "appear on a '%2' construct " + "with a '%3' clause%select{ with more than 1 argument|}0">; def err_acc_reduction_type : Error<"OpenACC 'reduction' variable must be of scalar type, sub-array, or a " "composite of scalar types;%select{| sub-array base}1 type is %0">; @@ -12779,13 +12779,14 @@ def err_acc_clause_in_clause_region def err_acc_gang_reduction_conflict : Error<"%select{OpenACC 'gang' clause with a 'dim' value greater than " "1|OpenACC 'reduction' clause}0 cannot " - "appear on the same 'loop' construct as a %select{'reduction' " + "appear on the same '%1' construct as a %select{'reduction' " "clause|'gang' clause with a 'dim' value greater than 1}0">; def err_acc_gang_reduction_numgangs_conflict - : Error<"OpenACC '%0' clause cannot appear on the same 'loop' construct " - "as a '%1' clause inside a compute construct with a " + : Error<"OpenACC '%0' clause cannot appear on the same '%2' construct as a " + "'%1' clause %select{inside a compute construct with a|and a}3 " "'num_gangs' clause with more than one argument">; -def err_reduction_op_mismatch + + def err_reduction_op_mismatch : Error<"OpenACC 'reduction' variable must have the same operator in all " "nested constructs (%0 vs %1)">; def err_acc_loop_variable_type diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h index a4132534686a81..170a6f4885c965 100644 --- a/clang/include/clang/Sema/SemaOpenACC.h +++ b/clang/include/clang/Sema/SemaOpenACC.h @@ -717,7 +717,8 @@ class SemaOpenACC : public SemaBase { // Does the checking for a 'gang' clause that needs to be done in dependent // and not dependent cases. OpenACCClause * - CheckGangClause(ArrayRef<const OpenACCClause *> ExistingClauses, + CheckGangClause(OpenACCDirectiveKind DirKind, + ArrayRef<const OpenACCClause *> ExistingClauses, SourceLocation BeginLoc, SourceLocation LParenLoc, ArrayRef<OpenACCGangKind> GangKinds, ArrayRef<Expr *> IntExprs, SourceLocation EndLoc); diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp index 18d58a43d265cf..62c3e778ab178d 100644 --- a/clang/lib/Sema/SemaOpenACC.cpp +++ b/clang/lib/Sema/SemaOpenACC.cpp @@ -719,11 +719,35 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitNumGangsClause( << /*NoArgs=*/1 << Clause.getDirectiveKind() << MaxArgs << Clause.getIntExprs().size(); + // OpenACC 3.3 Section 2.9.11: A reduction clause may not appear on a loop + // directive that has a gang clause and is within a compute construct that has + // a num_gangs clause with more than one explicit argument. + if (Clause.getIntExprs().size() > 1 && + isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind())) { + auto *GangClauseItr = + llvm::find_if(ExistingClauses, llvm::IsaPred<OpenACCGangClause>); + auto *ReductionClauseItr = + llvm::find_if(ExistingClauses, llvm::IsaPred<OpenACCReductionClause>); + + if (GangClauseItr != ExistingClauses.end() && + ReductionClauseItr != ExistingClauses.end()) { + SemaRef.Diag(Clause.getBeginLoc(), + diag::err_acc_gang_reduction_numgangs_conflict) + << OpenACCClauseKind::Reduction << OpenACCClauseKind::Gang + << Clause.getDirectiveKind() << /*is on combined directive=*/1; + SemaRef.Diag((*ReductionClauseItr)->getBeginLoc(), + diag::note_acc_previous_clause_here); + SemaRef.Diag((*GangClauseItr)->getBeginLoc(), + diag::note_acc_previous_clause_here); + return nullptr; + } + } + // OpenACC 3.3 Section 2.5.4: // A reduction clause may not appear on a parallel construct with a // num_gangs clause that has more than one argument. - // TODO: OpenACC: Reduction on Combined Construct needs to do this too. - if (Clause.getDirectiveKind() == OpenACCDirectiveKind::Parallel && + if ((Clause.getDirectiveKind() == OpenACCDirectiveKind::Parallel || + Clause.getDirectiveKind() == OpenACCDirectiveKind::ParallelLoop) && Clause.getIntExprs().size() > 1) { auto *Parallel = llvm::find_if(ExistingClauses, llvm::IsaPred<OpenACCReductionClause>); @@ -731,7 +755,8 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitNumGangsClause( if (Parallel != ExistingClauses.end()) { SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_reduction_num_gangs_conflict) - << Clause.getIntExprs().size(); + << /*>1 arg in first loc=*/1 << Clause.getClauseKind() + << Clause.getDirectiveKind() << OpenACCClauseKind::Reduction; SemaRef.Diag((*Parallel)->getBeginLoc(), diag::note_acc_previous_clause_here); return nullptr; @@ -739,7 +764,7 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitNumGangsClause( } // OpenACC 3.3 Section 2.9.2: - // An argument with no keyword or with the 'num' wkeyword is allowed only when + // An argument with no keyword or with the 'num' keyword is allowed only when // the 'num_gangs' does not appear on the 'kernel' construct. if (Clause.getDirectiveKind() == OpenACCDirectiveKind::KernelsLoop) { auto GangClauses = llvm::make_filter_range( @@ -1457,32 +1482,36 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitGangClause( // OpenACC 3.3 Section 2.9.11: A reduction clause may not appear on a loop // directive that has a gang clause and is within a compute construct that has // a num_gangs clause with more than one explicit argument. - // TODO OpenACC: When we implement reduction on combined constructs, we need - // to do this too. - if (Clause.getDirectiveKind() == OpenACCDirectiveKind::Loop && - SemaRef.getActiveComputeConstructInfo().Kind != - OpenACCDirectiveKind::Invalid) { + if ((Clause.getDirectiveKind() == OpenACCDirectiveKind::Loop && + SemaRef.getActiveComputeConstructInfo().Kind != + OpenACCDirectiveKind::Invalid) || + isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind())) { // num_gangs clause on the active compute construct. - auto *NumGangsClauseItr = - llvm::find_if(SemaRef.getActiveComputeConstructInfo().Clauses, - llvm::IsaPred<OpenACCNumGangsClause>); - - auto *ReductionClauseItr = - llvm::find_if(ExistingClauses, llvm::IsaPred<OpenACCReductionClause>); - - if (ReductionClauseItr != ExistingClauses.end() && - NumGangsClauseItr != - SemaRef.getActiveComputeConstructInfo().Clauses.end() && + auto ActiveComputeConstructContainer = + isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind()) + ? ExistingClauses + : SemaRef.getActiveComputeConstructInfo().Clauses; + auto *NumGangsClauseItr = llvm::find_if( + ActiveComputeConstructContainer, llvm::IsaPred<OpenACCNumGangsClause>); + + if (NumGangsClauseItr != ActiveComputeConstructContainer.end() && cast<OpenACCNumGangsClause>(*NumGangsClauseItr)->getIntExprs().size() > 1) { - SemaRef.Diag(Clause.getBeginLoc(), - diag::err_acc_gang_reduction_numgangs_conflict) - << OpenACCClauseKind::Gang << OpenACCClauseKind::Reduction; - SemaRef.Diag((*ReductionClauseItr)->getBeginLoc(), - diag::note_acc_previous_clause_here); - SemaRef.Diag((*NumGangsClauseItr)->getBeginLoc(), - diag::note_acc_previous_clause_here); - return nullptr; + auto *ReductionClauseItr = + llvm::find_if(ExistingClauses, llvm::IsaPred<OpenACCReductionClause>); + + if (ReductionClauseItr != ExistingClauses.end()) { + SemaRef.Diag(Clause.getBeginLoc(), + diag::err_acc_gang_reduction_numgangs_conflict) + << OpenACCClauseKind::Gang << OpenACCClauseKind::Reduction + << Clause.getDirectiveKind() + << isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind()); + SemaRef.Diag((*ReductionClauseItr)->getBeginLoc(), + diag::note_acc_previous_clause_here); + SemaRef.Diag((*NumGangsClauseItr)->getBeginLoc(), + diag::note_acc_previous_clause_here); + return nullptr; + } } } @@ -1563,9 +1592,9 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitGangClause( } } - return SemaRef.CheckGangClause(ExistingClauses, Clause.getBeginLoc(), - Clause.getLParenLoc(), GangKinds, IntExprs, - Clause.getEndLoc()); + return SemaRef.CheckGangClause(Clause.getDirectiveKind(), ExistingClauses, + Clause.getBeginLoc(), Clause.getLParenLoc(), + GangKinds, IntExprs, Clause.getEndLoc()); } OpenACCClause *SemaOpenACCClauseVisitor::VisitSeqClause( @@ -1609,41 +1638,39 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitSeqClause( OpenACCClause *SemaOpenACCClauseVisitor::VisitReductionClause( SemaOpenACC::OpenACCParsedClause &Clause) { - // Restrictions only properly implemented on 'compute' constructs, and - // 'compute' constructs are the only construct that can do anything with - // this yet, so skip/treat as unimplemented in this case. - // TODO: OpenACC: Remove check once we get combined constructs for this clause. - if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()) && - Clause.getDirectiveKind() != OpenACCDirectiveKind::Loop) - return isNotImplemented(); - // OpenACC 3.3 Section 2.9.11: A reduction clause may not appear on a loop // directive that has a gang clause and is within a compute construct that has // a num_gangs clause with more than one explicit argument. - if (Clause.getDirectiveKind() == OpenACCDirectiveKind::Loop && - SemaRef.getActiveComputeConstructInfo().Kind != - OpenACCDirectiveKind::Invalid) { + if ((Clause.getDirectiveKind() == OpenACCDirectiveKind::Loop && + SemaRef.getActiveComputeConstructInfo().Kind != + OpenACCDirectiveKind::Invalid) || + isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind())) { // num_gangs clause on the active compute construct. - auto *NumGangsClauseItr = - llvm::find_if(SemaRef.getActiveComputeConstructInfo().Clauses, - llvm::IsaPred<OpenACCNumGangsClause>); - - auto *GangClauseItr = - llvm::find_if(ExistingClauses, llvm::IsaPred<OpenACCGangClause>); - - if (GangClauseItr != ExistingClauses.end() && - NumGangsClauseItr != - SemaRef.getActiveComputeConstructInfo().Clauses.end() && + auto ActiveComputeConstructContainer = + isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind()) + ? ExistingClauses + : SemaRef.getActiveComputeConstructInfo().Clauses; + auto *NumGangsClauseItr = llvm::find_if( + ActiveComputeConstructContainer, llvm::IsaPred<OpenACCNumGangsClause>); + + if (NumGangsClauseItr != ActiveComputeConstructContainer.end() && cast<OpenACCNumGangsClause>(*NumGangsClauseItr)->getIntExprs().size() > 1) { - SemaRef.Diag(Clause.getBeginLoc(), - diag::err_acc_gang_reduction_numgangs_conflict) - << OpenACCClauseKind::Reduction << OpenACCClauseKind::Gang; - SemaRef.Diag((*GangClauseItr)->getBeginLoc(), - diag::note_acc_previous_clause_here); - SemaRef.Diag((*NumGangsClauseItr)->getBeginLoc(), - diag::note_acc_previous_clause_here); - return nullptr; + auto *GangClauseItr = + llvm::find_if(ExistingClauses, llvm::IsaPred<OpenACCGangClause>); + + if (GangClauseItr != ExistingClauses.end()) { + SemaRef.Diag(Clause.getBeginLoc(), + diag::err_acc_gang_reduction_numgangs_conflict) + << OpenACCClauseKind::Reduction << OpenACCClauseKind::Gang + << Clause.getDirectiveKind() + << isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind()); + SemaRef.Diag((*GangClauseItr)->getBeginLoc(), + diag::note_acc_previous_clause_here); + SemaRef.Diag((*NumGangsClauseItr)->getBeginLoc(), + diag::note_acc_previous_clause_here); + return nullptr; + } } } @@ -1667,7 +1694,8 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitReductionClause( // OpenACC 3.3 Section 2.5.4: // A reduction clause may not appear on a parallel construct with a // num_gangs clause that has more than one argument. - if (Clause.getDirectiveKind() == OpenACCDirectiveKind::Parallel) { + if (Clause.getDirectiveKind() == OpenACCDirectiveKind::Parallel || + Clause.getDirectiveKind() == OpenACCDirectiveKind::ParallelLoop) { auto NumGangsClauses = llvm::make_filter_range( ExistingClauses, llvm::IsaPred<OpenACCNumGangsClause>); @@ -1678,7 +1706,8 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitReductionClause( if (NumExprs > 1) { SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_reduction_num_gangs_conflict) - << NumExprs; + << /*>1 arg in first loc=*/0 << Clause.getClauseKind() + << Clause.getDirectiveKind() << OpenACCClauseKind::NumGangs; SemaRef.Diag(NGC->getBeginLoc(), diag::note_acc_previous_clause_here); return nullptr; } @@ -2624,7 +2653,8 @@ SemaOpenACC::CheckGangExpr(ArrayRef<const OpenACCClause *> ExistingClauses, } OpenACCClause * -SemaOpenACC::CheckGangClause(ArrayRef<const OpenACCClause *> ExistingClauses, +SemaOpenACC::CheckGangClause(OpenACCDirectiveKind DirKind, + ArrayRef<const OpenACCClause *> ExistingClauses, SourceLocation BeginLoc, SourceLocation LParenLoc, ArrayRef<OpenACCGangKind> GangKinds, ArrayRef<Expr *> IntExprs, SourceLocation EndLoc) { @@ -2649,7 +2679,7 @@ SemaOpenACC::CheckGangClause(ArrayRef<const OpenACCClause *> ExistingClauses, if (const auto *DimVal = dyn_cast<ConstantExpr>(DimExpr); DimVal && DimVal->getResultAsAPSInt() > 1) { Diag(DimVal->getBeginLoc(), diag::err_acc_gang_reduction_conflict) - << /*gang/reduction=*/0; + << /*gang/reduction=*/0 << DirKind; Diag((*ReductionItr)->getBeginLoc(), diag::note_acc_previous_clause_here); return nullptr; @@ -2666,30 +2696,29 @@ OpenACCClause *SemaOpenACC::CheckReductionClause( OpenACCDirectiveKind DirectiveKind, SourceLocation BeginLoc, SourceLocation LParenLoc, OpenACCReductionOperator ReductionOp, ArrayRef<Expr *> Vars, SourceLocation EndLoc) { - if (DirectiveKind == OpenACCDirectiveKind::Loop) { + if (DirectiveKind == OpenACCDirectiveKind::Loop || + isOpenACCCombinedDirectiveKind(DirectiveKind)) { // OpenACC 3.3 2.9.11: A reduction clause may not appear on a loop directive // that has a gang clause with a dim: argument whose value is greater // than 1. - const auto *GangItr = - llvm::find_if(ExistingClauses, llvm::IsaPred<OpenACCGangClause>); + const auto GangClauses = llvm::make_filter_range( + ExistingClauses, llvm::IsaPred<OpenACCGangClause>); - while (GangItr != ExistingClauses.end()) { - auto *GangClause = cast<OpenACCGangClause>(*GangItr); + for (auto *GC : GangClauses) { + const auto *GangClause = cast<OpenACCGangClause>(GC); for (unsigned I = 0; I < GangClause->getNumExprs(); ++I) { std::pair<OpenACCGangKind, const Expr *> EPair = GangClause->getExpr(I); - // We know there is only 1 on this gang, so move onto the next gang. if (EPair.first != OpenACCGangKind::Dim) - break; + continue; if (const auto *DimVal = dyn_cast<ConstantExpr>(EPair.second); DimVal && DimVal->getResultAsAPSInt() > 1) { Diag(BeginLoc, diag::err_acc_gang_reduction_conflict) - << /*reduction/gang=*/1; - Diag((*GangItr)->getBeginLoc(), diag::note_acc_previous_clause_here); + << /*reduction/gang=*/1 << DirectiveKind; + Diag(GangClause->getBeginLoc(), diag::note_acc_previous_clause_here); return nullptr; } } - ++GangItr; } } diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h index 81e515e7cb2a9a..02d2fc018e3c35 100644 --- a/clang/lib/Sema/TreeTransform.h +++ b/clang/lib/Sema/TreeTransform.h @@ -12037,7 +12037,8 @@ void OpenACCClauseTransform<Derived>::VisitGangClause( } NewClause = Self.getSema().OpenACC().CheckGangClause( - ExistingClauses, ParsedClause.getBeginLoc(), ParsedClause.getLParenLoc(), + ParsedClause.getDirectiveKind(), ExistingClauses, + ParsedClause.getBeginLoc(), ParsedClause.getLParenLoc(), TransformedGangKinds, TransformedIntExprs, ParsedClause.getEndLoc()); } } // namespace diff --git a/clang/test/AST/ast-print-openacc-combined-construct.cpp b/clang/test/AST/ast-print-openacc-combined-construct.cpp index 1a11f036b07aed..25fa29cbbe04e0 100644 --- a/clang/test/AST/ast-print-openacc-combined-construct.cpp +++ b/clang/test/AST/ast-print-openacc-combined-construct.cpp @@ -386,4 +386,31 @@ void foo() { #pragma acc serial loop vector for(int i = 0;i<5;++i); +//CHECK: #pragma acc parallel loop reduction(+: iPtr) +#pragma acc parallel loop reduction(+: iPtr) + for(int i = 0;i<5;++i); +//CHECK: #pragma acc serial loop reduction(*: i) +#pragma acc serial loop reduction(*: i) + for(int i = 0;i<5;++i); +//CHECK: #pragma acc kernels loop reduction(max: SomeB) +#pragma acc kernels loop reduction(max: SomeB) + for(int i = 0;i<5;++i); +//CHECK: #pragma acc parallel loop reduction(min: iPtr) +#pragma acc parallel loop reduction(min: iPtr) + for(int i = 0;i<5;++i); +//CHECK: #pragma acc serial loop reduction(&: i) +#pragma acc serial loop reduction(&: i) + for(int i = 0;i<5;++i); +//CHECK: #pragma acc kernels loop reduction(|: SomeB) +#pragma acc kernels loop reduction(|: SomeB) + for(int i = 0;i<5;++i); +//CHECK: #pragma acc parallel loop reduction(^: iPtr) +#pragma acc parallel loop reduction(^: iPtr) + for(int i = 0;i<5;++i); +//CHECK: #pragma acc serial loop reduction(&&: i) +#pragma acc serial loop reduction(&&: i) + for(int i = 0;i<5;++i); +//CHECK: #pragma acc kernels loop reduction(||: SomeB) +#pragma acc kernels loop reduction(||: SomeB) + for(int i = 0;i<5;++i); } diff --git a/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c b/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c index 45eede2e30f1f9..16bdcc177bfde3 100644 --- a/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c +++ b/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c @@ -1,8 +1,5 @@ // RUN: %clang_cc1 %s -fopenacc -verify -// TODO: OpenACC: A number of the 'not yet implemented' diagnostics interfere -// with the diagnostics we want to make here, so as we implement these, we need -// to replace the errors we should have. void uses() { #pragma acc parallel loop auto for(unsigned i = 0; i < 5; ++i); @@ -124,7 +121,6 @@ void uses() { // expected-warning@+1{{OpenACC clause name 'present_or_create' is a deprecated clause name and is now an alias for 'create'}} #pragma acc parallel loop auto present_or_create(Var) for(unsigned i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'reduction' not yet implemented}} #pragma acc parallel loop auto reduction(+:Var) for(unsigned i = 0; i < 5; ++i); #pragma acc parallel loop auto collapse(1) @@ -242,7 +238,6 @@ void uses() { // expected-warning@+1{{OpenACC clause name 'present_or_create' is a deprecated clause name and is now an alias for 'create'}} #pragma acc parallel loop present_or_create(Var) auto for(unsigned i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'reduction' not yet implemented}} #pragma acc parallel loop reduction(+:Var) auto for(unsigned i = 0; i < 5; ++i); #pragma acc parallel loop collapse(1) auto @@ -361,7 +356,6 @@ void uses() { // expected-warning@+1{{OpenACC clause name 'present_or_create' is a deprecated clause name and is now an alias for 'create'}} #pragma acc parallel loop independent present_or_create(Var) for(unsigned i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'reduction' not yet implemented}} #pragma acc parallel loop independent reduction(+:Var) for(unsigned i = 0; i < 5; ++i); #pragma acc parallel loop independent collapse(1) @@ -479,7 +473,6 @@ void uses() { // expected-warning@+1{{OpenACC clause name 'present_or_create' is a deprecated clause name and is now an alias for 'create'}} #pragma acc parallel loop present_or_create(Var) independent for(unsigned i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'reduction' not yet implemented}} #pragma acc parallel loop reduction(+:Var) independent for(unsigned i = 0; i < 5; ++i); #pragma acc parallel loop collapse(1) independent @@ -606,7 +599,6 @@ void uses() { // expected-warning@+1{{OpenACC clause name 'present_or_create' is a deprecated clause name and is now an alias for 'create'}} #pragma acc parallel loop seq present_or_create(Var) for(unsigned i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'reduction' not yet implemented}} #pragma acc parallel loop seq reduction(+:Var) for(unsigned i = 0; i < 5; ++i); #pragma acc parallel loop seq collapse(1) @@ -730,7 +722,6 @@ void uses() { // expected-warning@+1{{OpenACC clause name 'present_or_create' is a deprecated clause name and is now an alias for 'create'}} #pragma acc parallel loop present_or_create(Var) seq for(unsigned i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'reduction' not yet implemented}} #pragma acc parallel loop reduction(+:Var) seq for(unsigned i = 0; i < 5; ++i); #pragma acc parallel loop collapse(1) seq diff --git a/clang/test/SemaOpenACC/combined-construct-reduction-ast.cpp b/clang/test/SemaOpenACC/combined-construct-reduction-ast.cpp new file mode 100644 index 00000000000000..502b7a2613e402 --- /dev/null +++ b/clang/test/SemaOpenACC/combined-construct-reduction-ast.cpp @@ -0,0 +1,129 @@ +// 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 + +void NormalFunc(int i, float f) { + // CHECK: FunctionDecl{{.*}}NormalFunc + // CHECK-NEXT: ParmVarDecl + // CHECK-NEXT: ParmVarDecl + // CHECK-NEXT: CompoundStmt + +#pragma acc parallel loop reduction(+: i) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: reduction clause Operator: + + // CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar{{.*}} 'i' 'int' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc serial loop reduction(*: f) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} serial loop + // CHECK-NEXT: reduction clause Operator: * + // CHECK-NEXT: DeclRefExpr{{.*}} 'float' lvalue ParmVar{{.*}} 'f' 'float' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + + +#pragma acc kernels loop reduction(max: i) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop + // CHECK-NEXT: reduction clause Operator: max + // CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar{{.*}} 'i' 'int' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + + } + +template<typename T> +void TemplFunc() { + // CHECK: FunctionTemplateDecl{{.*}}TemplFunc + // CHECK-NEXT: TemplateTypeParmDecl + + // Match the prototype: + // CHECK-NEXT: FunctionDecl{{.*}}TemplFunc + // CHECK-NEXT: CompoundStmt + + T t; + // CHECK-NEXT: DeclStmt + // CHECK-NEXT: VarDecl{{.*}} t 'T' + +#pragma acc parallel loop reduction(+: t) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: reduction clause Operator: + + // CHECK-NEXT: DeclRefExpr{{.*}} 'T' lvalue Var{{.*}} 't' 'T' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +#pragma acc serial loop reduction(*: T::SomeFloat) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} serial loop + // CHECK-NEXT: reduction clause Operator: * + // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue + // CHECK-NEXT: NestedNameSpecifier TypeSpec 'T' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + + typename T::IntTy i; + // CHECK-NEXT: DeclStmt + // CHECK-NEXT: VarDecl{{.*}} i 'typename T::IntTy' + +#pragma acc kernels loop reduction(max: i) + for(int i = 0; i < 5; ++i); + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop + // CHECK-NEXT: reduction clause Operator: max + // CHECK-NEXT: DeclRefExpr{{.*}} 'typename T::IntTy' lvalue Var{{.*}} 'i' 'typename T::IntTy' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + + // Match the instantiation: + + // CHECK: FunctionDecl{{.*}}TemplFunc 'void ()' implicit_instantiation + // CHECK-NEXT: TemplateArgument type 'InstTy' + // CHECK-NEXT: RecordType{{.*}} 'InstTy' + // CHECK-NEXT: CXXRecord{{.*}} 'InstTy' + // CHECK-NEXT: CompoundStmt + + // CHECK-NEXT: DeclStmt + // CHECK-NEXT: VarDecl{{.*}} used t 'InstTy' + // CHECK-NEXT: CXXConstructExpr{{.*}} 'InstTy' 'void () noexcept' + + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop + // CHECK-NEXT: reduction clause Operator: + + // CHECK-NEXT: DeclRefExpr{{.*}} 'InstTy' lvalue Var{{.*}} 't' 'InstTy' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} serial loop + // CHECK-NEXT: reduction clause Operator: * + // CHECK-NEXT: DeclRefExpr{{.*}} 'const float' lvalue Var{{.*}} 'SomeFloat' 'const float' + // CHECK-NEXT: NestedNameSpecifier TypeSpec 'InstTy' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + + // CHECK-NEXT: DeclStmt + // CHECK-NEXT: VarDecl{{.*}} i 'typename InstTy::IntTy':'int' + + // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop + // CHECK-NEXT: reduction clause Operator: max + // CHECK-NEXT: DeclRefExpr{{.*}} 'typename InstTy::IntTy':'int' lvalue Var{{.*}} 'i' 'typename InstTy::IntTy':'int' + // CHECK-NEXT: ForStmt + // CHECK: NullStmt + +} + +struct InstTy { + using IntTy = int; + static constexpr float SomeFloat = 5.0; +}; + +void Instantiate() { + TemplFunc<InstTy>(); +} +#endif // PCH_HELPER diff --git a/clang/test/SemaOpenACC/combined-construct-reduction-clause.cpp b/clang/test/SemaOpenACC/combined-construct-reduction-clause.cpp new file mode 100644 index 00000000000000..082f1ed67a86f2 --- /dev/null +++ b/clang/test/SemaOpenACC/combined-construct-reduction-clause.cpp @@ -0,0 +1,169 @@ +// RUN: %clang_cc1 %s -fopenacc -verify + +struct CompositeOfScalars { + int I; + float F; + short J; + char C; + double D; + _Complex float CF; + _Complex double CD; +}; + +struct CompositeHasComposite { + int I; + float F; + short J; + char C; + double D; + _Complex float CF; + _Complex double CD; + struct CompositeOfScalars COS; // #COS_FIELD +}; + + // All of the type checking is done for compute and loop constructs, so only check the basics + the parts that are combined specific. +void uses(unsigned Parm) { + struct CompositeOfScalars CoS; + struct CompositeHasComposite ChC; + int I; + float F; + int Array[5]; + + // legal on all 3 kinds of combined constructs +#pragma acc parallel loop reduction(+:Parm) + for(int i = 0; i < 5; ++i); + +#pragma acc serial loop reduction(&: CoS, I, F) + for(int i = 0; i < 5; ++i); + +#pragma acc kernels loop reduction(min: CoS, Array[I], Array[0:I]) + for(int i = 0; i < 5; ++i); + + // expected-error@+2{{OpenACC 'reduction' composite variable must not have non-scalar field}} + // expected-note@#COS_FIELD{{invalid field is here}} +#pragma acc parallel loop reduction(&: ChC) + for(int i = 0; i < 5; ++i); + +#pragma acc kernels loop reduction(+:Parm) num_gangs(I) + for(int i = 0; i < 5; ++i); + // expected-error@+2{{OpenACC 'num_gangs' clause with more than 1 argument may not appear on a 'parallel loop' construct with a 'reduction' clause}} + // expected-note@+1{{previous clause is here}} +#pragma acc parallel loop reduction(+:Parm) num_gangs(I, I) + for(int i = 0; i < 5; ++i); + +#pragma acc kernels loop num_gangs(I) reduction(+:Parm) + for(int i = 0; i < 5; ++i); + // expected-error@+2{{OpenACC 'reduction' clause may not appear on a 'parallel loop' construct with a 'num_gangs' clause with more than 1 argument}} + // expected-note@+1{{previous clause is here}} +#pragma acc parallel loop num_gangs(I, I) reduction(+:Parm) + for(int i = 0; i < 5; ++i); + + // Reduction cannot appear on a loop with a 'gang' of dim>1. +#pragma acc parallel loop gang(dim:1) reduction(+:Parm) + for(int i = 0; i < 5; ++i); + // expected-error@+2{{OpenACC 'reduction' clause cannot appear on the same 'parallel loop' construct as a 'gang' clause with a 'dim' value greater than 1}} + // expected-note@+1{{previous clause is here}} +#pragma acc parallel loop gang(dim:2) reduction(+:Parm) + for(int i = 0; i < 5; ++i); +#pragma acc parallel loop reduction(+:Parm) gang(dim:1) + for(int i = 0; i < 5; ++i); + // expected-error@+2{{OpenACC 'gang' clause with a 'dim' value greater than 1 cannot appear on the same 'parallel loop' construct as a 'reduction' clause}} + // expected-note@+1{{previous clause is here}} +#pragma acc parallel loop reduction(+:Parm) gang(dim:2) + for(int i = 0; i < 5; ++i); + + // Reduction cannot appear on a loop with a gang and a num_gangs with >1 + // explicit argument. +#pragma acc kernels loop num_gangs(I) reduction(+:Parm) gang + for(int i = 0; i < 5; ++i); +#pragma acc kernels loop num_gangs(I) gang reduction(+:Parm) + for(int i = 0; i < 5; ++i); +#pragma acc kernels loop reduction(+:Parm) num_gangs(I) gang + for(int i = 0; i < 5; ++i); +#pragma acc kernels loop reduction(+:Parm) gang num_gangs(I) + for(int i = 0; i < 5; ++i); +#pragma acc kernels loop gang num_gangs(I) reduction(+:Parm) + for(int i = 0; i < 5; ++i); +#pragma acc kernels loop gang reduction(+:Parm) num_gangs(I) + for(int i = 0; i < 5; ++i); + + // expected-error@+2{{OpenACC 'reduction' clause may not appear on a 'parallel loop' construct with a 'num_gangs' clause with more than 1 argument}} + // expected-note@+1{{previous clause is here}} +#pragma acc parallel loop num_gangs(I, I) reduction(+:Parm) gang + for(int i = 0; i < 5; ++i); + // expected-error@+3{{OpenACC 'reduction' clause cannot appear on the same 'parallel loop' construct as a 'gang' clause and a 'num_gangs' clause with more than one argument}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc parallel loop num_gangs(I, I) gang reduction(+:Parm) + for(int i = 0; i < 5; ++i); + // expected-error@+2{{OpenACC 'num_gangs' clause with more than 1 argument may not appear on a 'parallel loop' construct with a 'reduction' clause}} + // expected-note@+1{{previous clause is here}} +#pragma acc parallel loop reduction(+:Parm) num_gangs(I, I) gang + for(int i = 0; i < 5; ++i); + // expected-error@+3{{OpenACC 'reduction' clause cannot appear on the same 'parallel loop' construct as a 'gang' clause and a 'num_gangs' clause with more than one argument}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc parallel loop reduction(+:Parm) gang num_gangs(I, I) + for(int i = 0; i < 5; ++i); + // expected-error@+3{{OpenACC 'reduction' clause cannot appear on the same 'parallel loop' construct as a 'gang' clause and a 'num_gangs' clause with more than one argument}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc parallel loop gang num_gangs(I, I) reduction(+:Parm) + for(int i = 0; i < 5; ++i); + // expected-error@+3{{OpenACC 'reduction' clause cannot appear on the same 'parallel loop' construct as a 'gang' clause and a 'num_gangs' clause with more than one argument}} + // expected-note@+2{{previous clause is here}} + // expected-note@+1{{previous clause is here}} +#pragma acc parallel loop gang reduction(+:Parm) num_gangs(I, I) + for(int i = 0; i < 5; ++i); + +#pragma acc parallel loop num_gangs(I) reduction(+:Parm) gang + for(int i = 0; i < 5; ++i); +#pragma acc parallel loop num_gangs(I) gang reduction(+:Parm) + for(int i = 0; i < 5; ++i); +#pragma acc parallel loop reduction(+:Parm) num_gangs(I) gang + for(int i = 0; i < 5; ++i); +#pragma acc parallel loop reduction(+:Parm) gang num_gangs(I) + for(int i = 0; i < 5; ++i); +#pragma acc parallel loop gang num_gangs(I) reduction(+:Parm) + for(int i = 0; i < 5; ++i); +#pragma acc parallel loop gang reduction(+:Parm) num_gangs(I) + for(int i = 0; i < 5; ++i); + +#pragma acc parallel loop reduction(+:I) + for(int i = 0; i < 5; ++i) { + // expected-error@+2{{OpenACC 'reduction' variable must have the same operator in all nested constructs (& vs +)}} + // expected-note@-3{{previous clause is here}} +#pragma acc loop reduction(&:I) + for(int i = 0; i < 5; ++i); + } +#pragma acc parallel loop reduction(+:I) + for(int i = 0; i < 5; ++i) { + // expected-error@+2{{OpenACC 'reduction' variable must have the same operator in all nested constructs (& vs +)}} + // expected-note@-3{{previous clause is here}} +#pragma acc parallel reduction(&:I) + for(int i = 0; i < 5; ++i); + } + +#pragma acc parallel loop reduction(+:I) + for(int i = 0; i < 5; ++i) { + // expected-error@+2{{OpenACC 'reduction' variable must have the same operator in all nested constructs (& vs +)}} + // expected-note@-3{{previous clause is here}} +#pragma acc parallel loop reduction(&:I) + for(int i = 0; i < 5; ++i); + } +#pragma acc loop reduction(+:I) + for(int i = 0; i < 5; ++i) { + // expected-error@+2{{OpenACC 'reduction' variable must have the same operator in all nested constructs (& vs +)}} + // expected-note@-3{{previous clause is here}} +#pragma acc parallel loop reduction(&:I) + for(int i = 0; i < 5; ++i); + } + +#pragma acc parallel reduction(+:I) + for(int i = 0; i < 5; ++i) { + // expected-error@+2{{OpenACC 'reduction' variable must have the same operator in all nested constructs (& vs +)}} + // expected-note@-3{{previous clause is here}} +#pragma acc parallel loop reduction(&:I) + for(int i = 0; i < 5; ++i); + } +} diff --git a/clang/test/SemaOpenACC/compute-construct-reduction-clause.c b/clang/test/SemaOpenACC/compute-construct-reduction-clause.c index 80310f0e7afc6d..fcc4ca2655c20a 100644 --- a/clang/test/SemaOpenACC/compute-construct-reduction-clause.c +++ b/clang/test/SemaOpenACC/compute-construct-reduction-clause.c @@ -41,12 +41,12 @@ void uses(unsigned Parm) { #pragma acc parallel num_gangs(IVar) reduction(+:IVar) while (1); - // expected-error@+2{{OpenACC 'reduction' clause may not appear on a 'parallel' construct with a 'num_gangs' clause with more than 1 argument, have 2}} + // expected-error@+2{{OpenACC 'num_gangs' clause with more than 1 argument may not appear on a 'parallel' construct with a 'reduction' clause}} // expected-note@+1{{previous clause is here}} #pragma acc parallel reduction(+:Parm) num_gangs(Parm, IVar) while (1); - // expected-error@+2{{OpenACC 'reduction' clause may not appear on a 'parallel' construct with a 'num_gangs' clause with more than 1 argument, have 2}} + // expected-error@+2{{OpenACC 'reduction' clause may not appear on a 'parallel' construct with a 'num_gangs' clause with more than 1 argument}} // expected-note@+1{{previous clause is here}} #pragma acc parallel num_gangs(Parm, IVar) reduction(+:Var) while (1); diff --git a/clang/test/SemaOpenACC/compute-construct-reduction-clause.cpp b/clang/test/SemaOpenACC/compute-construct-reduction-clause.cpp index 532dbb23871652..7372f683e2eb3e 100644 --- a/clang/test/SemaOpenACC/compute-construct-reduction-clause.cpp +++ b/clang/test/SemaOpenACC/compute-construct-reduction-clause.cpp @@ -41,12 +41,12 @@ void uses(unsigned Parm) { #pragma acc parallel num_gangs(IVar) reduction(+:Var) while (1); - // expected-error@+2{{OpenACC 'reduction' clause may not appear on a 'parallel' construct with a 'num_gangs' clause with more than 1 argument, have 2}} + // expected-error@+2{{OpenACC 'num_gangs' clause with more than 1 argument may not appear on a 'parallel' construct with a 'reduction' clause}} // expected-note@+1{{previous clause is here}} #pragma acc parallel reduction(+:Parm) num_gangs(Parm, IVar) while (1); - // expected-error@+2{{OpenACC 'reduction' clause may not appear on a 'parallel' construct with a 'num_gangs' clause with more than 1 argument, have 2}} + // expected-error@+2{{OpenACC 'reduction' clause may not appear on a 'parallel' construct with a 'num_gangs' clause with more than 1 argument}} // expected-note@+1{{previous clause is here}} #pragma acc parallel num_gangs(Parm, IVar) reduction(+:Var) while (1); @@ -116,12 +116,12 @@ void TemplUses(T Parm, U CoS, V ChC) { #pragma acc parallel num_gangs(Var) reduction(+:Var) while (1); - // expected-error@+2{{OpenACC 'reduction' clause may not appear on a 'parallel' construct with a 'num_gangs' clause with more than 1 argument, have 2}} + // expected-error@+2{{OpenACC 'num_gangs' clause with more than 1 argument may not appear on a 'parallel' construct with a 'reduction' clause}} // expected-note@+1{{previous clause is here}} #pragma acc parallel reduction(+:Parm) num_gangs(Parm, Var) while (1); - // expected-error@+2{{OpenACC 'reduction' clause may not appear on a 'parallel' construct with a 'num_gangs' clause with more than 1 argument, have 2}} + // expected-error@+2{{OpenACC 'reduction' clause may not appear on a 'parallel' construct with a 'num_gangs' clause with more than 1 argument}} // expected-note@+1{{previous clause is here}} #pragma acc parallel num_gangs(Parm, Var) reduction(+:Var) while (1); _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
