[clang] 7d89ebf - [OpenACC] Implement 'reduction' for combined constructs.
via cfe-commits
cfe-commits at lists.llvm.org
Mon Dec 9 14:06:49 PST 2024
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 at +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 at +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 at +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 at +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 at +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 at +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 at +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 at +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 at +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 at +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 at +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 at +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 at +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 at +2{{OpenACC 'num_gangs' clause with more than 1 argument may not appear on a 'parallel loop' construct with a 'reduction' clause}}
+ // expected-note at +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 at +2{{OpenACC 'reduction' clause may not appear on a 'parallel loop' construct with a 'num_gangs' clause with more than 1 argument}}
+ // expected-note at +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 at +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 at +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 at +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 at +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 at +2{{OpenACC 'reduction' clause may not appear on a 'parallel loop' construct with a 'num_gangs' clause with more than 1 argument}}
+ // expected-note at +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 at +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 at +2{{previous clause is here}}
+ // expected-note at +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 at +2{{OpenACC 'num_gangs' clause with more than 1 argument may not appear on a 'parallel loop' construct with a 'reduction' clause}}
+ // expected-note at +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 at +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 at +2{{previous clause is here}}
+ // expected-note at +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 at +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 at +2{{previous clause is here}}
+ // expected-note at +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 at +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 at +2{{previous clause is here}}
+ // expected-note at +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 at +2{{OpenACC 'reduction' variable must have the same operator in all nested constructs (& vs +)}}
+ // expected-note at -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 at +2{{OpenACC 'reduction' variable must have the same operator in all nested constructs (& vs +)}}
+ // expected-note at -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 at +2{{OpenACC 'reduction' variable must have the same operator in all nested constructs (& vs +)}}
+ // expected-note at -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 at +2{{OpenACC 'reduction' variable must have the same operator in all nested constructs (& vs +)}}
+ // expected-note at -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 at +2{{OpenACC 'reduction' variable must have the same operator in all nested constructs (& vs +)}}
+ // expected-note at -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 at +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 at +2{{OpenACC 'num_gangs' clause with more than 1 argument may not appear on a 'parallel' construct with a 'reduction' clause}}
// expected-note at +1{{previous clause is here}}
#pragma acc parallel reduction(+:Parm) num_gangs(Parm, IVar)
while (1);
- // expected-error at +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 at +2{{OpenACC 'reduction' clause may not appear on a 'parallel' construct with a 'num_gangs' clause with more than 1 argument}}
// expected-note at +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 at +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 at +2{{OpenACC 'num_gangs' clause with more than 1 argument may not appear on a 'parallel' construct with a 'reduction' clause}}
// expected-note at +1{{previous clause is here}}
#pragma acc parallel reduction(+:Parm) num_gangs(Parm, IVar)
while (1);
- // expected-error at +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 at +2{{OpenACC 'reduction' clause may not appear on a 'parallel' construct with a 'num_gangs' clause with more than 1 argument}}
// expected-note at +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 at +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 at +2{{OpenACC 'num_gangs' clause with more than 1 argument may not appear on a 'parallel' construct with a 'reduction' clause}}
// expected-note at +1{{previous clause is here}}
#pragma acc parallel reduction(+:Parm) num_gangs(Parm, Var)
while (1);
- // expected-error at +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 at +2{{OpenACC 'reduction' clause may not appear on a 'parallel' construct with a 'num_gangs' clause with more than 1 argument}}
// expected-note at +1{{previous clause is here}}
#pragma acc parallel num_gangs(Parm, Var) reduction(+:Var)
while (1);
More information about the cfe-commits
mailing list