Skip to content

Commit 7d89ebf

Browse files
erichkeaneerichkeane
committedDec 9, 2024·
[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.
·
llvmorg-21-initllvmorg-20.1.0-rc1
1 parent 44cd8f0 commit 7d89ebf

10 files changed

+445
-97
lines changed
 

‎clang/include/clang/Basic/DiagnosticSemaKinds.td

+8-7
Original file line numberDiff line numberDiff line change
@@ -12716,9 +12716,9 @@ def err_acc_clause_cannot_combine
1271612716
: Error<"OpenACC clause '%0' may not appear on the same construct as a "
1271712717
"'%1' clause on a '%2' construct">;
1271812718
def err_acc_reduction_num_gangs_conflict
12719-
: Error<
12720-
"OpenACC 'reduction' clause may not appear on a 'parallel' construct "
12721-
"with a 'num_gangs' clause with more than 1 argument, have %0">;
12719+
: Error<"OpenACC '%1' clause %select{|with more than 1 argument }0may not "
12720+
"appear on a '%2' construct "
12721+
"with a '%3' clause%select{ with more than 1 argument|}0">;
1272212722
def err_acc_reduction_type
1272312723
: Error<"OpenACC 'reduction' variable must be of scalar type, sub-array, or a "
1272412724
"composite of scalar types;%select{| sub-array base}1 type is %0">;
@@ -12779,13 +12779,14 @@ def err_acc_clause_in_clause_region
1277912779
def err_acc_gang_reduction_conflict
1278012780
: Error<"%select{OpenACC 'gang' clause with a 'dim' value greater than "
1278112781
"1|OpenACC 'reduction' clause}0 cannot "
12782-
"appear on the same 'loop' construct as a %select{'reduction' "
12782+
"appear on the same '%1' construct as a %select{'reduction' "
1278312783
"clause|'gang' clause with a 'dim' value greater than 1}0">;
1278412784
def err_acc_gang_reduction_numgangs_conflict
12785-
: Error<"OpenACC '%0' clause cannot appear on the same 'loop' construct "
12786-
"as a '%1' clause inside a compute construct with a "
12785+
: Error<"OpenACC '%0' clause cannot appear on the same '%2' construct as a "
12786+
"'%1' clause %select{inside a compute construct with a|and a}3 "
1278712787
"'num_gangs' clause with more than one argument">;
12788-
def err_reduction_op_mismatch
12788+
12789+
def err_reduction_op_mismatch
1278912790
: Error<"OpenACC 'reduction' variable must have the same operator in all "
1279012791
"nested constructs (%0 vs %1)">;
1279112792
def err_acc_loop_variable_type

‎clang/include/clang/Sema/SemaOpenACC.h

+2-1
Original file line numberDiff line numberDiff line change
@@ -717,7 +717,8 @@ class SemaOpenACC : public SemaBase {
717717
// Does the checking for a 'gang' clause that needs to be done in dependent
718718
// and not dependent cases.
719719
OpenACCClause *
720-
CheckGangClause(ArrayRef<const OpenACCClause *> ExistingClauses,
720+
CheckGangClause(OpenACCDirectiveKind DirKind,
721+
ArrayRef<const OpenACCClause *> ExistingClauses,
721722
SourceLocation BeginLoc, SourceLocation LParenLoc,
722723
ArrayRef<OpenACCGangKind> GangKinds,
723724
ArrayRef<Expr *> IntExprs, SourceLocation EndLoc);

‎clang/lib/Sema/SemaOpenACC.cpp

+102-73
Original file line numberDiff line numberDiff line change
@@ -719,27 +719,52 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitNumGangsClause(
719719
<< /*NoArgs=*/1 << Clause.getDirectiveKind() << MaxArgs
720720
<< Clause.getIntExprs().size();
721721

722+
// OpenACC 3.3 Section 2.9.11: A reduction clause may not appear on a loop
723+
// directive that has a gang clause and is within a compute construct that has
724+
// a num_gangs clause with more than one explicit argument.
725+
if (Clause.getIntExprs().size() > 1 &&
726+
isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind())) {
727+
auto *GangClauseItr =
728+
llvm::find_if(ExistingClauses, llvm::IsaPred<OpenACCGangClause>);
729+
auto *ReductionClauseItr =
730+
llvm::find_if(ExistingClauses, llvm::IsaPred<OpenACCReductionClause>);
731+
732+
if (GangClauseItr != ExistingClauses.end() &&
733+
ReductionClauseItr != ExistingClauses.end()) {
734+
SemaRef.Diag(Clause.getBeginLoc(),
735+
diag::err_acc_gang_reduction_numgangs_conflict)
736+
<< OpenACCClauseKind::Reduction << OpenACCClauseKind::Gang
737+
<< Clause.getDirectiveKind() << /*is on combined directive=*/1;
738+
SemaRef.Diag((*ReductionClauseItr)->getBeginLoc(),
739+
diag::note_acc_previous_clause_here);
740+
SemaRef.Diag((*GangClauseItr)->getBeginLoc(),
741+
diag::note_acc_previous_clause_here);
742+
return nullptr;
743+
}
744+
}
745+
722746
// OpenACC 3.3 Section 2.5.4:
723747
// A reduction clause may not appear on a parallel construct with a
724748
// num_gangs clause that has more than one argument.
725-
// TODO: OpenACC: Reduction on Combined Construct needs to do this too.
726-
if (Clause.getDirectiveKind() == OpenACCDirectiveKind::Parallel &&
749+
if ((Clause.getDirectiveKind() == OpenACCDirectiveKind::Parallel ||
750+
Clause.getDirectiveKind() == OpenACCDirectiveKind::ParallelLoop) &&
727751
Clause.getIntExprs().size() > 1) {
728752
auto *Parallel =
729753
llvm::find_if(ExistingClauses, llvm::IsaPred<OpenACCReductionClause>);
730754

731755
if (Parallel != ExistingClauses.end()) {
732756
SemaRef.Diag(Clause.getBeginLoc(),
733757
diag::err_acc_reduction_num_gangs_conflict)
734-
<< Clause.getIntExprs().size();
758+
<< /*>1 arg in first loc=*/1 << Clause.getClauseKind()
759+
<< Clause.getDirectiveKind() << OpenACCClauseKind::Reduction;
735760
SemaRef.Diag((*Parallel)->getBeginLoc(),
736761
diag::note_acc_previous_clause_here);
737762
return nullptr;
738763
}
739764
}
740765

741766
// OpenACC 3.3 Section 2.9.2:
742-
// An argument with no keyword or with the 'num' wkeyword is allowed only when
767+
// An argument with no keyword or with the 'num' keyword is allowed only when
743768
// the 'num_gangs' does not appear on the 'kernel' construct.
744769
if (Clause.getDirectiveKind() == OpenACCDirectiveKind::KernelsLoop) {
745770
auto GangClauses = llvm::make_filter_range(
@@ -1457,32 +1482,36 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitGangClause(
14571482
// OpenACC 3.3 Section 2.9.11: A reduction clause may not appear on a loop
14581483
// directive that has a gang clause and is within a compute construct that has
14591484
// a num_gangs clause with more than one explicit argument.
1460-
// TODO OpenACC: When we implement reduction on combined constructs, we need
1461-
// to do this too.
1462-
if (Clause.getDirectiveKind() == OpenACCDirectiveKind::Loop &&
1463-
SemaRef.getActiveComputeConstructInfo().Kind !=
1464-
OpenACCDirectiveKind::Invalid) {
1485+
if ((Clause.getDirectiveKind() == OpenACCDirectiveKind::Loop &&
1486+
SemaRef.getActiveComputeConstructInfo().Kind !=
1487+
OpenACCDirectiveKind::Invalid) ||
1488+
isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind())) {
14651489
// num_gangs clause on the active compute construct.
1466-
auto *NumGangsClauseItr =
1467-
llvm::find_if(SemaRef.getActiveComputeConstructInfo().Clauses,
1468-
llvm::IsaPred<OpenACCNumGangsClause>);
1469-
1470-
auto *ReductionClauseItr =
1471-
llvm::find_if(ExistingClauses, llvm::IsaPred<OpenACCReductionClause>);
1472-
1473-
if (ReductionClauseItr != ExistingClauses.end() &&
1474-
NumGangsClauseItr !=
1475-
SemaRef.getActiveComputeConstructInfo().Clauses.end() &&
1490+
auto ActiveComputeConstructContainer =
1491+
isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind())
1492+
? ExistingClauses
1493+
: SemaRef.getActiveComputeConstructInfo().Clauses;
1494+
auto *NumGangsClauseItr = llvm::find_if(
1495+
ActiveComputeConstructContainer, llvm::IsaPred<OpenACCNumGangsClause>);
1496+
1497+
if (NumGangsClauseItr != ActiveComputeConstructContainer.end() &&
14761498
cast<OpenACCNumGangsClause>(*NumGangsClauseItr)->getIntExprs().size() >
14771499
1) {
1478-
SemaRef.Diag(Clause.getBeginLoc(),
1479-
diag::err_acc_gang_reduction_numgangs_conflict)
1480-
<< OpenACCClauseKind::Gang << OpenACCClauseKind::Reduction;
1481-
SemaRef.Diag((*ReductionClauseItr)->getBeginLoc(),
1482-
diag::note_acc_previous_clause_here);
1483-
SemaRef.Diag((*NumGangsClauseItr)->getBeginLoc(),
1484-
diag::note_acc_previous_clause_here);
1485-
return nullptr;
1500+
auto *ReductionClauseItr =
1501+
llvm::find_if(ExistingClauses, llvm::IsaPred<OpenACCReductionClause>);
1502+
1503+
if (ReductionClauseItr != ExistingClauses.end()) {
1504+
SemaRef.Diag(Clause.getBeginLoc(),
1505+
diag::err_acc_gang_reduction_numgangs_conflict)
1506+
<< OpenACCClauseKind::Gang << OpenACCClauseKind::Reduction
1507+
<< Clause.getDirectiveKind()
1508+
<< isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind());
1509+
SemaRef.Diag((*ReductionClauseItr)->getBeginLoc(),
1510+
diag::note_acc_previous_clause_here);
1511+
SemaRef.Diag((*NumGangsClauseItr)->getBeginLoc(),
1512+
diag::note_acc_previous_clause_here);
1513+
return nullptr;
1514+
}
14861515
}
14871516
}
14881517

@@ -1563,9 +1592,9 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitGangClause(
15631592
}
15641593
}
15651594

1566-
return SemaRef.CheckGangClause(ExistingClauses, Clause.getBeginLoc(),
1567-
Clause.getLParenLoc(), GangKinds, IntExprs,
1568-
Clause.getEndLoc());
1595+
return SemaRef.CheckGangClause(Clause.getDirectiveKind(), ExistingClauses,
1596+
Clause.getBeginLoc(), Clause.getLParenLoc(),
1597+
GangKinds, IntExprs, Clause.getEndLoc());
15691598
}
15701599

15711600
OpenACCClause *SemaOpenACCClauseVisitor::VisitSeqClause(
@@ -1609,41 +1638,39 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitSeqClause(
16091638

16101639
OpenACCClause *SemaOpenACCClauseVisitor::VisitReductionClause(
16111640
SemaOpenACC::OpenACCParsedClause &Clause) {
1612-
// Restrictions only properly implemented on 'compute' constructs, and
1613-
// 'compute' constructs are the only construct that can do anything with
1614-
// this yet, so skip/treat as unimplemented in this case.
1615-
// TODO: OpenACC: Remove check once we get combined constructs for this clause.
1616-
if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()) &&
1617-
Clause.getDirectiveKind() != OpenACCDirectiveKind::Loop)
1618-
return isNotImplemented();
1619-
16201641
// OpenACC 3.3 Section 2.9.11: A reduction clause may not appear on a loop
16211642
// directive that has a gang clause and is within a compute construct that has
16221643
// a num_gangs clause with more than one explicit argument.
1623-
if (Clause.getDirectiveKind() == OpenACCDirectiveKind::Loop &&
1624-
SemaRef.getActiveComputeConstructInfo().Kind !=
1625-
OpenACCDirectiveKind::Invalid) {
1644+
if ((Clause.getDirectiveKind() == OpenACCDirectiveKind::Loop &&
1645+
SemaRef.getActiveComputeConstructInfo().Kind !=
1646+
OpenACCDirectiveKind::Invalid) ||
1647+
isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind())) {
16261648
// num_gangs clause on the active compute construct.
1627-
auto *NumGangsClauseItr =
1628-
llvm::find_if(SemaRef.getActiveComputeConstructInfo().Clauses,
1629-
llvm::IsaPred<OpenACCNumGangsClause>);
1630-
1631-
auto *GangClauseItr =
1632-
llvm::find_if(ExistingClauses, llvm::IsaPred<OpenACCGangClause>);
1633-
1634-
if (GangClauseItr != ExistingClauses.end() &&
1635-
NumGangsClauseItr !=
1636-
SemaRef.getActiveComputeConstructInfo().Clauses.end() &&
1649+
auto ActiveComputeConstructContainer =
1650+
isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind())
1651+
? ExistingClauses
1652+
: SemaRef.getActiveComputeConstructInfo().Clauses;
1653+
auto *NumGangsClauseItr = llvm::find_if(
1654+
ActiveComputeConstructContainer, llvm::IsaPred<OpenACCNumGangsClause>);
1655+
1656+
if (NumGangsClauseItr != ActiveComputeConstructContainer.end() &&
16371657
cast<OpenACCNumGangsClause>(*NumGangsClauseItr)->getIntExprs().size() >
16381658
1) {
1639-
SemaRef.Diag(Clause.getBeginLoc(),
1640-
diag::err_acc_gang_reduction_numgangs_conflict)
1641-
<< OpenACCClauseKind::Reduction << OpenACCClauseKind::Gang;
1642-
SemaRef.Diag((*GangClauseItr)->getBeginLoc(),
1643-
diag::note_acc_previous_clause_here);
1644-
SemaRef.Diag((*NumGangsClauseItr)->getBeginLoc(),
1645-
diag::note_acc_previous_clause_here);
1646-
return nullptr;
1659+
auto *GangClauseItr =
1660+
llvm::find_if(ExistingClauses, llvm::IsaPred<OpenACCGangClause>);
1661+
1662+
if (GangClauseItr != ExistingClauses.end()) {
1663+
SemaRef.Diag(Clause.getBeginLoc(),
1664+
diag::err_acc_gang_reduction_numgangs_conflict)
1665+
<< OpenACCClauseKind::Reduction << OpenACCClauseKind::Gang
1666+
<< Clause.getDirectiveKind()
1667+
<< isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind());
1668+
SemaRef.Diag((*GangClauseItr)->getBeginLoc(),
1669+
diag::note_acc_previous_clause_here);
1670+
SemaRef.Diag((*NumGangsClauseItr)->getBeginLoc(),
1671+
diag::note_acc_previous_clause_here);
1672+
return nullptr;
1673+
}
16471674
}
16481675
}
16491676

@@ -1667,7 +1694,8 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitReductionClause(
16671694
// OpenACC 3.3 Section 2.5.4:
16681695
// A reduction clause may not appear on a parallel construct with a
16691696
// num_gangs clause that has more than one argument.
1670-
if (Clause.getDirectiveKind() == OpenACCDirectiveKind::Parallel) {
1697+
if (Clause.getDirectiveKind() == OpenACCDirectiveKind::Parallel ||
1698+
Clause.getDirectiveKind() == OpenACCDirectiveKind::ParallelLoop) {
16711699
auto NumGangsClauses = llvm::make_filter_range(
16721700
ExistingClauses, llvm::IsaPred<OpenACCNumGangsClause>);
16731701

@@ -1678,7 +1706,8 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitReductionClause(
16781706
if (NumExprs > 1) {
16791707
SemaRef.Diag(Clause.getBeginLoc(),
16801708
diag::err_acc_reduction_num_gangs_conflict)
1681-
<< NumExprs;
1709+
<< /*>1 arg in first loc=*/0 << Clause.getClauseKind()
1710+
<< Clause.getDirectiveKind() << OpenACCClauseKind::NumGangs;
16821711
SemaRef.Diag(NGC->getBeginLoc(), diag::note_acc_previous_clause_here);
16831712
return nullptr;
16841713
}
@@ -2624,7 +2653,8 @@ SemaOpenACC::CheckGangExpr(ArrayRef<const OpenACCClause *> ExistingClauses,
26242653
}
26252654

26262655
OpenACCClause *
2627-
SemaOpenACC::CheckGangClause(ArrayRef<const OpenACCClause *> ExistingClauses,
2656+
SemaOpenACC::CheckGangClause(OpenACCDirectiveKind DirKind,
2657+
ArrayRef<const OpenACCClause *> ExistingClauses,
26282658
SourceLocation BeginLoc, SourceLocation LParenLoc,
26292659
ArrayRef<OpenACCGangKind> GangKinds,
26302660
ArrayRef<Expr *> IntExprs, SourceLocation EndLoc) {
@@ -2649,7 +2679,7 @@ SemaOpenACC::CheckGangClause(ArrayRef<const OpenACCClause *> ExistingClauses,
26492679
if (const auto *DimVal = dyn_cast<ConstantExpr>(DimExpr);
26502680
DimVal && DimVal->getResultAsAPSInt() > 1) {
26512681
Diag(DimVal->getBeginLoc(), diag::err_acc_gang_reduction_conflict)
2652-
<< /*gang/reduction=*/0;
2682+
<< /*gang/reduction=*/0 << DirKind;
26532683
Diag((*ReductionItr)->getBeginLoc(),
26542684
diag::note_acc_previous_clause_here);
26552685
return nullptr;
@@ -2666,30 +2696,29 @@ OpenACCClause *SemaOpenACC::CheckReductionClause(
26662696
OpenACCDirectiveKind DirectiveKind, SourceLocation BeginLoc,
26672697
SourceLocation LParenLoc, OpenACCReductionOperator ReductionOp,
26682698
ArrayRef<Expr *> Vars, SourceLocation EndLoc) {
2669-
if (DirectiveKind == OpenACCDirectiveKind::Loop) {
2699+
if (DirectiveKind == OpenACCDirectiveKind::Loop ||
2700+
isOpenACCCombinedDirectiveKind(DirectiveKind)) {
26702701
// OpenACC 3.3 2.9.11: A reduction clause may not appear on a loop directive
26712702
// that has a gang clause with a dim: argument whose value is greater
26722703
// than 1.
2673-
const auto *GangItr =
2674-
llvm::find_if(ExistingClauses, llvm::IsaPred<OpenACCGangClause>);
2704+
const auto GangClauses = llvm::make_filter_range(
2705+
ExistingClauses, llvm::IsaPred<OpenACCGangClause>);
26752706

2676-
while (GangItr != ExistingClauses.end()) {
2677-
auto *GangClause = cast<OpenACCGangClause>(*GangItr);
2707+
for (auto *GC : GangClauses) {
2708+
const auto *GangClause = cast<OpenACCGangClause>(GC);
26782709
for (unsigned I = 0; I < GangClause->getNumExprs(); ++I) {
26792710
std::pair<OpenACCGangKind, const Expr *> EPair = GangClause->getExpr(I);
2680-
// We know there is only 1 on this gang, so move onto the next gang.
26812711
if (EPair.first != OpenACCGangKind::Dim)
2682-
break;
2712+
continue;
26832713

26842714
if (const auto *DimVal = dyn_cast<ConstantExpr>(EPair.second);
26852715
DimVal && DimVal->getResultAsAPSInt() > 1) {
26862716
Diag(BeginLoc, diag::err_acc_gang_reduction_conflict)
2687-
<< /*reduction/gang=*/1;
2688-
Diag((*GangItr)->getBeginLoc(), diag::note_acc_previous_clause_here);
2717+
<< /*reduction/gang=*/1 << DirectiveKind;
2718+
Diag(GangClause->getBeginLoc(), diag::note_acc_previous_clause_here);
26892719
return nullptr;
26902720
}
26912721
}
2692-
++GangItr;
26932722
}
26942723
}
26952724

‎clang/lib/Sema/TreeTransform.h

+2-1
Original file line numberDiff line numberDiff line change
@@ -12037,7 +12037,8 @@ void OpenACCClauseTransform<Derived>::VisitGangClause(
1203712037
}
1203812038

1203912039
NewClause = Self.getSema().OpenACC().CheckGangClause(
12040-
ExistingClauses, ParsedClause.getBeginLoc(), ParsedClause.getLParenLoc(),
12040+
ParsedClause.getDirectiveKind(), ExistingClauses,
12041+
ParsedClause.getBeginLoc(), ParsedClause.getLParenLoc(),
1204112042
TransformedGangKinds, TransformedIntExprs, ParsedClause.getEndLoc());
1204212043
}
1204312044
} // namespace

‎clang/test/AST/ast-print-openacc-combined-construct.cpp

+27
Original file line numberDiff line numberDiff line change
@@ -386,4 +386,31 @@ void foo() {
386386
#pragma acc serial loop vector
387387
for(int i = 0;i<5;++i);
388388

389+
//CHECK: #pragma acc parallel loop reduction(+: iPtr)
390+
#pragma acc parallel loop reduction(+: iPtr)
391+
for(int i = 0;i<5;++i);
392+
//CHECK: #pragma acc serial loop reduction(*: i)
393+
#pragma acc serial loop reduction(*: i)
394+
for(int i = 0;i<5;++i);
395+
//CHECK: #pragma acc kernels loop reduction(max: SomeB)
396+
#pragma acc kernels loop reduction(max: SomeB)
397+
for(int i = 0;i<5;++i);
398+
//CHECK: #pragma acc parallel loop reduction(min: iPtr)
399+
#pragma acc parallel loop reduction(min: iPtr)
400+
for(int i = 0;i<5;++i);
401+
//CHECK: #pragma acc serial loop reduction(&: i)
402+
#pragma acc serial loop reduction(&: i)
403+
for(int i = 0;i<5;++i);
404+
//CHECK: #pragma acc kernels loop reduction(|: SomeB)
405+
#pragma acc kernels loop reduction(|: SomeB)
406+
for(int i = 0;i<5;++i);
407+
//CHECK: #pragma acc parallel loop reduction(^: iPtr)
408+
#pragma acc parallel loop reduction(^: iPtr)
409+
for(int i = 0;i<5;++i);
410+
//CHECK: #pragma acc serial loop reduction(&&: i)
411+
#pragma acc serial loop reduction(&&: i)
412+
for(int i = 0;i<5;++i);
413+
//CHECK: #pragma acc kernels loop reduction(||: SomeB)
414+
#pragma acc kernels loop reduction(||: SomeB)
415+
for(int i = 0;i<5;++i);
389416
}

0 commit comments

Comments
 (0)
Please sign in to comment.