diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index f185a98720f5..0a245e2077f6 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 a4132534686a..170a6f4885c9 100644 --- a/clang/include/clang/Sema/SemaOpenACC.h +++ b/clang/include/clang/Sema/SemaOpenACC.h @@ -717,7 +717,8 @@ public: // Does the checking for a 'gang' clause that needs to be done in dependent // and not dependent cases. OpenACCClause * - CheckGangClause(ArrayRef ExistingClauses, + CheckGangClause(OpenACCDirectiveKind DirKind, + ArrayRef ExistingClauses, SourceLocation BeginLoc, SourceLocation LParenLoc, ArrayRef GangKinds, ArrayRef IntExprs, SourceLocation EndLoc); diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp index 18d58a43d265..62c3e778ab17 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); + auto *ReductionClauseItr = + llvm::find_if(ExistingClauses, llvm::IsaPred); + + 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); @@ -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); + auto ActiveComputeConstructContainer = + isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind()) + ? ExistingClauses + : SemaRef.getActiveComputeConstructInfo().Clauses; + auto *NumGangsClauseItr = llvm::find_if( + ActiveComputeConstructContainer, llvm::IsaPred); - auto *ReductionClauseItr = - llvm::find_if(ExistingClauses, llvm::IsaPred); - - if (ReductionClauseItr != ExistingClauses.end() && - NumGangsClauseItr != - SemaRef.getActiveComputeConstructInfo().Clauses.end() && + if (NumGangsClauseItr != ActiveComputeConstructContainer.end() && cast(*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); + + 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); + auto ActiveComputeConstructContainer = + isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind()) + ? ExistingClauses + : SemaRef.getActiveComputeConstructInfo().Clauses; + auto *NumGangsClauseItr = llvm::find_if( + ActiveComputeConstructContainer, llvm::IsaPred); - auto *GangClauseItr = - llvm::find_if(ExistingClauses, llvm::IsaPred); - - if (GangClauseItr != ExistingClauses.end() && - NumGangsClauseItr != - SemaRef.getActiveComputeConstructInfo().Clauses.end() && + if (NumGangsClauseItr != ActiveComputeConstructContainer.end() && cast(*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); + + 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); @@ -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 ExistingClauses, } OpenACCClause * -SemaOpenACC::CheckGangClause(ArrayRef ExistingClauses, +SemaOpenACC::CheckGangClause(OpenACCDirectiveKind DirKind, + ArrayRef ExistingClauses, SourceLocation BeginLoc, SourceLocation LParenLoc, ArrayRef GangKinds, ArrayRef IntExprs, SourceLocation EndLoc) { @@ -2649,7 +2679,7 @@ SemaOpenACC::CheckGangClause(ArrayRef ExistingClauses, if (const auto *DimVal = dyn_cast(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 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); + const auto GangClauses = llvm::make_filter_range( + ExistingClauses, llvm::IsaPred); - while (GangItr != ExistingClauses.end()) { - auto *GangClause = cast(*GangItr); + for (auto *GC : GangClauses) { + const auto *GangClause = cast(GC); for (unsigned I = 0; I < GangClause->getNumExprs(); ++I) { std::pair 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(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 81e515e7cb2a..02d2fc018e3c 100644 --- a/clang/lib/Sema/TreeTransform.h +++ b/clang/lib/Sema/TreeTransform.h @@ -12037,7 +12037,8 @@ void OpenACCClauseTransform::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 1a11f036b07a..25fa29cbbe04 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 45eede2e30f1..16bdcc177bfd 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 000000000000..502b7a2613e4 --- /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 +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{{.*}} '' 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(); +} +#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 000000000000..082f1ed67a86 --- /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 80310f0e7afc..fcc4ca2655c2 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 532dbb238716..7372f683e2eb 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);