[clang] 0b0b153 - [OpenACC] Implement 'reduction' Sema on 'loop' construct
via cfe-commits
cfe-commits at lists.llvm.org
Fri Nov 1 08:53:01 PDT 2024
Author: erichkeane
Date: 2024-11-01T08:52:15-07:00
New Revision: 0b0b15309435df19a93fefbc556c57a83da0d3f6
URL: https://github.com/llvm/llvm-project/commit/0b0b15309435df19a93fefbc556c57a83da0d3f6
DIFF: https://github.com/llvm/llvm-project/commit/0b0b15309435df19a93fefbc556c57a83da0d3f6.diff
LOG: [OpenACC] Implement 'reduction' Sema on 'loop' construct
The reduction clause has some minor restrictions on the variable
references that are implementable, so this implements those. Others
require reachability analysis, so this patch documents that we're not
going to do that in the CFE(or at least save it for the MLIR passes).
Added:
clang/test/SemaOpenACC/loop-construct-reduction-ast.cpp
clang/test/SemaOpenACC/loop-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-loop-construct.cpp
clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
Removed:
################################################################################
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 34ff49d7238a7f..d697e6d61afa9a 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12718,7 +12718,18 @@ def err_acc_num_arg_conflict
def err_acc_clause_in_clause_region
: Error<"loop with a '%0' clause may not exist in the region of a '%1' "
"clause%select{| on a 'kernels' compute construct}2">;
-
+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' "
+ "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 "
+ "'num_gangs' clause with more than one argument">;
+def err_reduction_op_mismatch
+ : Error<"OpenACC 'reduction' variable must have the same operator in all "
+ "nested constructs (%0 vs %1)">;
// AMDGCN builtins diagnostics
def err_amdgcn_global_load_lds_size_invalid_value : Error<"invalid size value">;
def note_amdgcn_global_load_lds_size_valid_value : Note<"size must be 1, 2, or 4">;
diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h
index a34551f6b5a192..cc5b0c4ece1ebe 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -68,6 +68,7 @@ class SemaOpenACC : public SemaBase {
/// used to diagnose if there are multiple 'for' loops at any one level.
LLVM_PREFERRED_TYPE(bool)
unsigned CurLevelHasLoopAlready : 1;
+
} LoopInfo{/*TopLevelLoopSeen=*/false, /*CurLevelHasLoopAlready=*/false};
/// The 'collapse' clause requires quite a bit of checking while
@@ -109,6 +110,14 @@ class SemaOpenACC : public SemaBase {
bool TileDepthSatisfied = true;
} TileInfo;
+ /// A list of the active reduction clauses, which allows us to check that all
+ /// vars on nested constructs for the same reduction var have the same
+ /// reduction operator. Currently this is enforced against all constructs
+ /// despite the rule being in the 'loop' section. By current reading, this
+ /// should apply to all anyway, but we may need to make this more like the
+ /// 'loop' clause enforcement, where this is 'blocked' by a compute construct.
+ llvm::SmallVector<OpenACCReductionClause *> ActiveReductionClauses;
+
public:
ComputeConstructInfo &getActiveComputeConstructInfo() {
return ActiveComputeConstructInfo;
@@ -615,7 +624,9 @@ class SemaOpenACC : public SemaBase {
/// Called while semantically analyzing the reduction clause, ensuring the var
/// is the correct kind of reference.
- ExprResult CheckReductionVar(Expr *VarExpr);
+ ExprResult CheckReductionVar(OpenACCDirectiveKind DirectiveKind,
+ OpenACCReductionOperator ReductionOp,
+ Expr *VarExpr);
/// Called to check the 'var' type is a variable of pointer type, necessary
/// for 'deviceptr' and 'attach' clauses. Returns true on success.
@@ -634,6 +645,22 @@ class SemaOpenACC : public SemaBase {
// Check a single expression on a gang clause.
ExprResult CheckGangExpr(OpenACCGangKind GK, Expr *E);
+ // Does the checking for a 'gang' clause that needs to be done in dependent
+ // and not dependent cases.
+ OpenACCClause *
+ CheckGangClause(ArrayRef<const OpenACCClause *> ExistingClauses,
+ SourceLocation BeginLoc, SourceLocation LParenLoc,
+ ArrayRef<OpenACCGangKind> GangKinds,
+ ArrayRef<Expr *> IntExprs, SourceLocation EndLoc);
+ // Does the checking for a 'reduction ' clause that needs to be done in
+ // dependent and not dependent cases.
+ OpenACCClause *
+ CheckReductionClause(ArrayRef<const OpenACCClause *> ExistingClauses,
+ OpenACCDirectiveKind DirectiveKind,
+ SourceLocation BeginLoc, SourceLocation LParenLoc,
+ OpenACCReductionOperator ReductionOp,
+ ArrayRef<Expr *> Vars, SourceLocation EndLoc);
+
ExprResult BuildOpenACCAsteriskSizeExpr(SourceLocation AsteriskLoc);
ExprResult ActOnOpenACCAsteriskSizeExpr(SourceLocation AsteriskLoc);
@@ -686,6 +713,7 @@ class SemaOpenACC : public SemaBase {
SourceLocation OldLoopWorkerClauseLoc;
SourceLocation OldLoopVectorClauseLoc;
llvm::SmallVector<OpenACCLoopConstruct *> ParentlessLoopConstructs;
+ llvm::SmallVector<OpenACCReductionClause *> ActiveReductionClauses;
LoopInConstructRAII LoopRAII;
public:
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index aadcd7183e713b..3073e87af90fa7 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -1224,6 +1224,36 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitGangClause(
if (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) {
+ // 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() &&
+ 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;
+ }
+ }
+
llvm::SmallVector<OpenACCGangKind> GangKinds;
llvm::SmallVector<Expr *> IntExprs;
@@ -1316,9 +1346,9 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitGangClause(
return nullptr;
}
- return OpenACCGangClause::Create(Ctx, Clause.getBeginLoc(),
- Clause.getLParenLoc(), GangKinds, IntExprs,
- Clause.getEndLoc());
+ return SemaRef.CheckGangClause(ExistingClauses, Clause.getBeginLoc(),
+ Clause.getLParenLoc(), GangKinds, IntExprs,
+ Clause.getEndLoc());
}
OpenACCClause *SemaOpenACCClauseVisitor::VisitSeqClause(
@@ -1367,9 +1397,57 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitReductionClause(
// 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.
- if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()))
+ 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) {
+ // 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() &&
+ 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;
+ }
+ }
+
+ // OpenACC3.3 Section 2.9.11: If a variable is involved in a reduction that
+ // spans multiple nested loops where two or more of those loops have
+ // associated loop directives, a reduction clause containing that variable
+ // must appear on each of those loop directives.
+ //
+ // This can't really be implemented in the CFE, as this requires a level of
+ // rechability/useage analysis that we're not really wanting to get into.
+ // Additionally, I'm alerted that this restriction is one that the middle-end
+ // can just 'figure out' as an extension and isn't really necessary.
+ //
+ // OpenACC3.3 Section 2.9.11: Every 'var' in a reduction clause appearing on
+ // an orphaned loop construct must be private.
+ //
+ // This again is something we cannot really diagnose, as it requires we see
+ // all the uses/scopes of all variables referenced. The middle end/MLIR might
+ // be able to diagnose this.
+
// 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.
@@ -1394,15 +1472,17 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitReductionClause(
SmallVector<Expr *> ValidVars;
for (Expr *Var : Clause.getVarList()) {
- ExprResult Res = SemaRef.CheckReductionVar(Var);
+ ExprResult Res = SemaRef.CheckReductionVar(Clause.getDirectiveKind(),
+ Clause.getReductionOp(), Var);
if (Res.isUsable())
ValidVars.push_back(Res.get());
}
- return OpenACCReductionClause::Create(
- Ctx, Clause.getBeginLoc(), Clause.getLParenLoc(), Clause.getReductionOp(),
- ValidVars, Clause.getEndLoc());
+ return SemaRef.CheckReductionClause(
+ ExistingClauses, Clause.getDirectiveKind(), Clause.getBeginLoc(),
+ Clause.getLParenLoc(), Clause.getReductionOp(), ValidVars,
+ Clause.getEndLoc());
}
OpenACCClause *SemaOpenACCClauseVisitor::VisitCollapseClause(
@@ -1425,6 +1505,16 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitCollapseClause(
LoopCount.get(), Clause.getEndLoc());
}
+void CollectActiveReductionClauses(
+ llvm::SmallVector<OpenACCReductionClause *> &ActiveClauses,
+ ArrayRef<OpenACCClause *> CurClauses) {
+ for (auto *CurClause : CurClauses) {
+ if (auto *RedClause = dyn_cast<OpenACCReductionClause>(CurClause);
+ RedClause && !RedClause->getVarList().empty())
+ ActiveClauses.push_back(RedClause);
+ }
+}
+
} // namespace
SemaOpenACC::SemaOpenACC(Sema &S) : SemaBase(S) {}
@@ -1437,11 +1527,14 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
DirKind(DK), OldLoopGangClauseOnKernelLoc(S.LoopGangClauseOnKernelLoc),
OldLoopWorkerClauseLoc(S.LoopWorkerClauseLoc),
OldLoopVectorClauseLoc(S.LoopVectorClauseLoc),
+ ActiveReductionClauses(S.ActiveReductionClauses),
LoopRAII(SemaRef, /*PreserveDepth=*/false) {
+
// Compute constructs end up taking their 'loop'.
if (DirKind == OpenACCDirectiveKind::Parallel ||
DirKind == OpenACCDirectiveKind::Serial ||
DirKind == OpenACCDirectiveKind::Kernels) {
+ CollectActiveReductionClauses(S.ActiveReductionClauses, Clauses);
SemaRef.ActiveComputeConstructInfo.Kind = DirKind;
SemaRef.ActiveComputeConstructInfo.Clauses = Clauses;
SemaRef.ParentlessLoopConstructs.swap(ParentlessLoopConstructs);
@@ -1456,6 +1549,7 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
SemaRef.LoopWorkerClauseLoc = {};
SemaRef.LoopVectorClauseLoc = {};
} else if (DirKind == OpenACCDirectiveKind::Loop) {
+ CollectActiveReductionClauses(S.ActiveReductionClauses, Clauses);
SetCollapseInfoBeforeAssociatedStmt(UnInstClauses, Clauses);
SetTileInfoBeforeAssociatedStmt(UnInstClauses, Clauses);
@@ -1559,6 +1653,7 @@ SemaOpenACC::AssociatedStmtRAII::~AssociatedStmtRAII() {
SemaRef.LoopGangClauseOnKernelLoc = OldLoopGangClauseOnKernelLoc;
SemaRef.LoopWorkerClauseLoc = OldLoopWorkerClauseLoc;
SemaRef.LoopVectorClauseLoc = OldLoopVectorClauseLoc;
+ SemaRef.ActiveReductionClauses.swap(ActiveReductionClauses);
if (DirKind == OpenACCDirectiveKind::Parallel ||
DirKind == OpenACCDirectiveKind::Serial ||
@@ -1610,6 +1705,68 @@ SemaOpenACC::ActOnClause(ArrayRef<const OpenACCClause *> ExistingClauses,
}
+namespace {
+// Return true if the two vars refer to the same variable, for the purposes of
+// equality checking.
+bool areVarsEqual(Expr *VarExpr1, Expr *VarExpr2) {
+ if (VarExpr1->isInstantiationDependent() ||
+ VarExpr2->isInstantiationDependent())
+ return false;
+
+ VarExpr1 = VarExpr1->IgnoreParenCasts();
+ VarExpr2 = VarExpr2->IgnoreParenCasts();
+
+ // Legal expressions can be: Scalar variable reference, sub-array, array
+ // element, or composite variable member.
+
+ // Sub-array.
+ if (isa<ArraySectionExpr>(VarExpr1)) {
+ auto *Expr2AS = dyn_cast<ArraySectionExpr>(VarExpr2);
+ if (!Expr2AS)
+ return false;
+
+ auto *Expr1AS = cast<ArraySectionExpr>(VarExpr1);
+
+ if (!areVarsEqual(Expr1AS->getBase(), Expr2AS->getBase()))
+ return false;
+ // We could possibly check to see if the ranges aren't overlapping, but it
+ // isn't clear that the rules allow this.
+ return true;
+ }
+
+ // Array-element.
+ if (isa<ArraySubscriptExpr>(VarExpr1)) {
+ auto *Expr2AS = dyn_cast<ArraySubscriptExpr>(VarExpr2);
+ if (!Expr2AS)
+ return false;
+
+ auto *Expr1AS = cast<ArraySubscriptExpr>(VarExpr1);
+
+ if (!areVarsEqual(Expr1AS->getBase(), Expr2AS->getBase()))
+ return false;
+
+ // We could possibly check to see if the elements referenced aren't the
+ // same, but it isn't clear by reading of the standard that this is allowed
+ // (and that the 'var' refered to isn't the array).
+ return true;
+ }
+
+ // Scalar variable reference, or composite variable.
+ if (isa<DeclRefExpr>(VarExpr1)) {
+ auto *Expr2DRE = dyn_cast<DeclRefExpr>(VarExpr2);
+ if (!Expr2DRE)
+ return false;
+
+ auto *Expr1DRE = cast<DeclRefExpr>(VarExpr1);
+
+ return Expr1DRE->getDecl()->getMostRecentDecl() ==
+ Expr2DRE->getDecl()->getMostRecentDecl();
+ }
+
+ llvm_unreachable("Unknown variable type encountered");
+}
+} // namespace
+
/// OpenACC 3.3 section 2.5.15:
/// At a mininmum, the supported data types include ... the numerical data types
/// in C, C++, and Fortran.
@@ -1617,7 +1774,9 @@ SemaOpenACC::ActOnClause(ArrayRef<const OpenACCClause *> ExistingClauses,
/// If the reduction var is a composite variable, each
/// member of the composite variable must be a supported datatype for the
/// reduction operation.
-ExprResult SemaOpenACC::CheckReductionVar(Expr *VarExpr) {
+ExprResult SemaOpenACC::CheckReductionVar(OpenACCDirectiveKind DirectiveKind,
+ OpenACCReductionOperator ReductionOp,
+ Expr *VarExpr) {
VarExpr = VarExpr->IgnoreParenCasts();
auto TypeIsValid = [](QualType Ty) {
@@ -1667,6 +1826,28 @@ ExprResult SemaOpenACC::CheckReductionVar(Expr *VarExpr) {
return ExprError();
}
+ // OpenACC3.3: 2.9.11: Reduction clauses on nested constructs for the same
+ // reduction 'var' must have the same reduction operator.
+ if (!VarExpr->isInstantiationDependent()) {
+
+ for (const OpenACCReductionClause *RClause : ActiveReductionClauses) {
+ if (RClause->getReductionOp() == ReductionOp)
+ break;
+
+ for (Expr *OldVarExpr : RClause->getVarList()) {
+ if (OldVarExpr->isInstantiationDependent())
+ continue;
+
+ if (areVarsEqual(VarExpr, OldVarExpr)) {
+ Diag(VarExpr->getExprLoc(), diag::err_reduction_op_mismatch)
+ << ReductionOp << RClause->getReductionOp();
+ Diag(OldVarExpr->getExprLoc(), diag::note_acc_previous_clause_here);
+ return ExprError();
+ }
+ }
+ }
+ }
+
return VarExpr;
}
@@ -2223,6 +2404,81 @@ ExprResult SemaOpenACC::CheckGangExpr(OpenACCGangKind GK, Expr *E) {
llvm_unreachable("Compute construct directive not handled?");
}
+OpenACCClause *
+SemaOpenACC::CheckGangClause(ArrayRef<const OpenACCClause *> ExistingClauses,
+ SourceLocation BeginLoc, SourceLocation LParenLoc,
+ ArrayRef<OpenACCGangKind> GangKinds,
+ ArrayRef<Expr *> IntExprs, SourceLocation EndLoc) {
+ // 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 *ReductionItr =
+ llvm::find_if(ExistingClauses, llvm::IsaPred<OpenACCReductionClause>);
+
+ if (ReductionItr != ExistingClauses.end()) {
+ const auto GangZip = llvm::zip_equal(GangKinds, IntExprs);
+ const auto GangItr = llvm::find_if(GangZip, [](const auto &Tuple) {
+ return std::get<0>(Tuple) == OpenACCGangKind::Dim;
+ });
+
+ if (GangItr != GangZip.end()) {
+ const Expr *DimExpr = std::get<1>(*GangItr);
+
+ assert(
+ (DimExpr->isInstantiationDependent() || isa<ConstantExpr>(DimExpr)) &&
+ "Improperly formed gang argument");
+ if (const auto *DimVal = dyn_cast<ConstantExpr>(DimExpr);
+ DimVal && DimVal->getResultAsAPSInt() > 1) {
+ Diag(DimVal->getBeginLoc(), diag::err_acc_gang_reduction_conflict)
+ << /*gang/reduction=*/0;
+ Diag((*ReductionItr)->getBeginLoc(),
+ diag::note_acc_previous_clause_here);
+ return nullptr;
+ }
+ }
+ }
+
+ return OpenACCGangClause::Create(getASTContext(), BeginLoc, LParenLoc,
+ GangKinds, IntExprs, EndLoc);
+}
+
+OpenACCClause *SemaOpenACC::CheckReductionClause(
+ ArrayRef<const OpenACCClause *> ExistingClauses,
+ OpenACCDirectiveKind DirectiveKind, SourceLocation BeginLoc,
+ SourceLocation LParenLoc, OpenACCReductionOperator ReductionOp,
+ ArrayRef<Expr *> Vars, SourceLocation EndLoc) {
+ if (DirectiveKind == OpenACCDirectiveKind::Loop) {
+ // 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>);
+
+ while (GangItr != ExistingClauses.end()) {
+ auto *GangClause = cast<OpenACCGangClause>(*GangItr);
+ 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;
+
+ 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);
+ return nullptr;
+ }
+ }
+ ++GangItr;
+ }
+ }
+
+ auto *Ret = OpenACCReductionClause::Create(
+ getASTContext(), BeginLoc, LParenLoc, ReductionOp, Vars, EndLoc);
+ return Ret;
+}
+
ExprResult SemaOpenACC::CheckTileSizeExpr(Expr *SizeExpr) {
if (!SizeExpr)
return ExprError();
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index aa34300dfa58db..15ba022b096ac3 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -11935,15 +11935,16 @@ void OpenACCClauseTransform<Derived>::VisitReductionClause(
SmallVector<Expr *> ValidVars;
for (Expr *Var : TransformedVars) {
- ExprResult Res = Self.getSema().OpenACC().CheckReductionVar(Var);
+ ExprResult Res = Self.getSema().OpenACC().CheckReductionVar(
+ ParsedClause.getDirectiveKind(), C.getReductionOp(), Var);
if (Res.isUsable())
ValidVars.push_back(Res.get());
}
- NewClause = OpenACCReductionClause::Create(
- Self.getSema().getASTContext(), ParsedClause.getBeginLoc(),
- ParsedClause.getLParenLoc(), C.getReductionOp(), ValidVars,
- ParsedClause.getEndLoc());
+ NewClause = Self.getSema().OpenACC().CheckReductionClause(
+ ExistingClauses, ParsedClause.getDirectiveKind(),
+ ParsedClause.getBeginLoc(), ParsedClause.getLParenLoc(),
+ C.getReductionOp(), ValidVars, ParsedClause.getEndLoc());
}
template <typename Derived>
@@ -12018,10 +12019,9 @@ void OpenACCClauseTransform<Derived>::VisitGangClause(
TransformedIntExprs.push_back(ER.get());
}
- NewClause = OpenACCGangClause::Create(
- Self.getSema().getASTContext(), ParsedClause.getBeginLoc(),
- ParsedClause.getLParenLoc(), TransformedGangKinds, TransformedIntExprs,
- ParsedClause.getEndLoc());
+ NewClause = Self.getSema().OpenACC().CheckGangClause(
+ ExistingClauses, ParsedClause.getBeginLoc(), ParsedClause.getLParenLoc(),
+ TransformedGangKinds, TransformedIntExprs, ParsedClause.getEndLoc());
}
} // namespace
template <typename Derived>
diff --git a/clang/test/AST/ast-print-openacc-loop-construct.cpp b/clang/test/AST/ast-print-openacc-loop-construct.cpp
index c0ca274f38dc2c..5338334af745f2 100644
--- a/clang/test/AST/ast-print-openacc-loop-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-loop-construct.cpp
@@ -291,4 +291,34 @@ void foo() {
#pragma acc loop vector
for(;;);
+ int *iPtr;
+ bool SomeB;
+
+//CHECK: #pragma acc loop reduction(+: iPtr)
+#pragma acc loop reduction(+: iPtr)
+ for(;;);
+//CHECK: #pragma acc loop reduction(*: i)
+#pragma acc loop reduction(*: i)
+ for(;;);
+//CHECK: #pragma acc loop reduction(max: SomeB)
+#pragma acc loop reduction(max: SomeB)
+ for(;;);
+//CHECK: #pragma acc loop reduction(min: iPtr)
+#pragma acc loop reduction(min: iPtr)
+ for(;;);
+//CHECK: #pragma acc loop reduction(&: i)
+#pragma acc loop reduction(&: i)
+ for(;;);
+//CHECK: #pragma acc loop reduction(|: SomeB)
+#pragma acc loop reduction(|: SomeB)
+ for(;;);
+//CHECK: #pragma acc loop reduction(^: iPtr)
+#pragma acc loop reduction(^: iPtr)
+ for(;;);
+//CHECK: #pragma acc loop reduction(&&: i)
+#pragma acc loop reduction(&&: i)
+ for(;;);
+//CHECK: #pragma acc loop reduction(||: SomeB)
+#pragma acc loop reduction(||: SomeB)
+ for(;;);
}
diff --git a/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c b/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
index ab10857e3cd858..e938c8f983055a 100644
--- a/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
+++ b/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
@@ -133,7 +133,6 @@ void uses() {
// expected-error at +1{{OpenACC 'present_or_create' clause is not valid on 'loop' directive}}
#pragma acc loop auto present_or_create(Var)
for(;;);
- // expected-warning at +1{{OpenACC clause 'reduction' not yet implemented}}
#pragma acc loop auto reduction(+:Var)
for(;;);
#pragma acc loop auto collapse(1)
@@ -268,7 +267,6 @@ void uses() {
// expected-error at +1{{OpenACC 'present_or_create' clause is not valid on 'loop' directive}}
#pragma acc loop present_or_create(Var) auto
for(;;);
- // expected-warning at +1{{OpenACC clause 'reduction' not yet implemented}}
#pragma acc loop reduction(+:Var) auto
for(;;);
#pragma acc loop collapse(1) auto
@@ -404,7 +402,6 @@ void uses() {
// expected-error at +1{{OpenACC 'present_or_create' clause is not valid on 'loop' directive}}
#pragma acc loop independent present_or_create(Var)
for(;;);
- // expected-warning at +1{{OpenACC clause 'reduction' not yet implemented}}
#pragma acc loop independent reduction(+:Var)
for(;;);
#pragma acc loop independent collapse(1)
@@ -539,7 +536,6 @@ void uses() {
// expected-error at +1{{OpenACC 'present_or_create' clause is not valid on 'loop' directive}}
#pragma acc loop present_or_create(Var) independent
for(;;);
- // expected-warning at +1{{OpenACC clause 'reduction' not yet implemented}}
#pragma acc loop reduction(+:Var) independent
for(;;);
#pragma acc loop collapse(1) independent
@@ -683,7 +679,6 @@ void uses() {
// expected-error at +1{{OpenACC 'present_or_create' clause is not valid on 'loop' directive}}
#pragma acc loop seq present_or_create(Var)
for(;;);
- // expected-warning at +1{{OpenACC clause 'reduction' not yet implemented}}
#pragma acc loop seq reduction(+:Var)
for(;;);
#pragma acc loop seq collapse(1)
@@ -824,7 +819,6 @@ void uses() {
// expected-error at +1{{OpenACC 'present_or_create' clause is not valid on 'loop' directive}}
#pragma acc loop present_or_create(Var) seq
for(;;);
- // expected-warning at +1{{OpenACC clause 'reduction' not yet implemented}}
#pragma acc loop reduction(+:Var) seq
for(;;);
#pragma acc loop collapse(1) seq
diff --git a/clang/test/SemaOpenACC/loop-construct-reduction-ast.cpp b/clang/test/SemaOpenACC/loop-construct-reduction-ast.cpp
new file mode 100644
index 00000000000000..d40f2672a4fea0
--- /dev/null
+++ b/clang/test/SemaOpenACC/loop-construct-reduction-ast.cpp
@@ -0,0 +1,375 @@
+// 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 loop reduction(+: i)
+ for(;;);
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: reduction clause Operator: +
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar{{.*}} 'i' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+
+#pragma acc loop reduction(*: f)
+ for(;;);
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: reduction clause Operator: *
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'float' lvalue ParmVar{{.*}} 'f' 'float'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+
+#pragma acc loop reduction(max: i)
+ for(;;);
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: reduction clause Operator: max
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar{{.*}} 'i' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+
+#pragma acc loop reduction(min: f)
+ for(;;);
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: reduction clause Operator: min
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'float' lvalue ParmVar{{.*}} 'f' 'float'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+
+#pragma acc loop reduction(&: i)
+ for(;;);
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: reduction clause Operator: &
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar{{.*}} 'i' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+
+#pragma acc loop reduction(|: f)
+ for(;;);
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: reduction clause Operator: |
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'float' lvalue ParmVar{{.*}} 'f' 'float'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+
+
+#pragma acc loop reduction(^: i)
+ for(;;);
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: reduction clause Operator: ^
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar{{.*}} 'i' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+
+#pragma acc loop reduction(&&: f)
+ for(;;);
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: reduction clause Operator: &&
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'float' lvalue ParmVar{{.*}} 'f' 'float'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+
+
+#pragma acc loop reduction(||: i)
+ for(;;);
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: reduction clause Operator: ||
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar{{.*}} 'i' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: 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 loop reduction(+: t)
+ for(;;);
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: reduction clause Operator: +
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'T' lvalue Var{{.*}} 't' 'T'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+
+#pragma acc loop reduction(*: T::SomeFloat)
+ for(;;);
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: reduction clause Operator: *
+ // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue
+ // CHECK-NEXT: NestedNameSpecifier TypeSpec 'T'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+
+ typename T::IntTy i;
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} i 'typename T::IntTy'
+
+#pragma acc loop reduction(max: i)
+ for(;;);
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: reduction clause Operator: max
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'typename T::IntTy' lvalue Var{{.*}} 'i' 'typename T::IntTy'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+
+#pragma acc loop reduction(min: t)
+ for(;;);
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: reduction clause Operator: min
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'T' lvalue Var{{.*}} 't' 'T'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+
+#pragma acc loop reduction(&: T::SomeFloat)
+ for(;;);
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: reduction clause Operator: &
+ // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue
+ // CHECK-NEXT: NestedNameSpecifier TypeSpec 'T'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+
+#pragma acc loop reduction(|: i)
+ for(;;);
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: reduction clause Operator: |
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'typename T::IntTy' lvalue Var{{.*}} 'i' 'typename T::IntTy'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+
+#pragma acc loop reduction(^: t)
+ for(;;);
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: reduction clause Operator: ^
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'T' lvalue Var{{.*}} 't' 'T'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+
+#pragma acc loop reduction(&&: T::SomeFloat)
+ for(;;);
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: reduction clause Operator: &&
+ // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue
+ // CHECK-NEXT: NestedNameSpecifier TypeSpec 'T'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+
+#pragma acc loop reduction(||: i)
+ for(;;);
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: reduction clause Operator: ||
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'typename T::IntTy' lvalue Var{{.*}} 'i' 'typename T::IntTy'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: 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: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: reduction clause Operator: +
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'InstTy' lvalue Var{{.*}} 't' 'InstTy'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+ //
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: reduction clause Operator: *
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'const float' lvalue Var{{.*}} 'SomeFloat' 'const float'
+ // CHECK-NEXT: NestedNameSpecifier TypeSpec 'InstTy'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+ //
+ // CHECK-NEXT: DeclStmt
+ // CHECK-NEXT: VarDecl{{.*}} i 'typename InstTy::IntTy':'int'
+ //
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: reduction clause Operator: max
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'typename InstTy::IntTy':'int' lvalue Var{{.*}} 'i' 'typename InstTy::IntTy':'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+ //
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: reduction clause Operator: min
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'InstTy' lvalue Var{{.*}} 't' 'InstTy'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+ //
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: reduction clause Operator: &
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'const float' lvalue Var{{.*}} 'SomeFloat' 'const float'
+ // CHECK-NEXT: NestedNameSpecifier TypeSpec 'InstTy'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+ //
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: reduction clause Operator: |
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'typename InstTy::IntTy':'int' lvalue Var{{.*}} 'i' 'typename InstTy::IntTy':'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+ //
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: reduction clause Operator: ^
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'InstTy' lvalue Var{{.*}} 't' 'InstTy'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+ //
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: reduction clause Operator: &&
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'const float' lvalue Var{{.*}} 'SomeFloat' 'const float'
+ // CHECK-NEXT: NestedNameSpecifier TypeSpec 'InstTy'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+ //
+ // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+ // CHECK-NEXT: reduction clause Operator: ||
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'typename InstTy::IntTy':'int' lvalue Var{{.*}} 'i' 'typename InstTy::IntTy':'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: <<<NULL>>>
+ // CHECK-NEXT: NullStmt
+
+}
+
+struct BoolConversion{ operator bool() const;};
+struct InstTy {
+ using IntTy = int;
+ static constexpr float SomeFloat = 5.0;
+ static constexpr BoolConversion BC;
+};
+
+void Instantiate() {
+ TemplFunc<InstTy>();
+}
+
+#endif // PCH_HELPER
diff --git a/clang/test/SemaOpenACC/loop-construct-reduction-clause.cpp b/clang/test/SemaOpenACC/loop-construct-reduction-clause.cpp
new file mode 100644
index 00000000000000..60151da5989c9e
--- /dev/null
+++ b/clang/test/SemaOpenACC/loop-construct-reduction-clause.cpp
@@ -0,0 +1,366 @@
+// 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
+};
+void uses() {
+
+ int I;
+ float F;
+ int Array[5];
+ CompositeOfScalars CoS;
+ CompositeHasComposite ChC;
+
+#pragma acc serial
+ {
+#pragma acc loop reduction(+:CoS, I, F)
+ for(;;){}
+ }
+
+#pragma acc serial
+ {
+ // expected-error at +1{{OpenACC 'reduction' variable must be of scalar type, sub-array, or a composite of scalar types; type is 'int[5]'}}
+#pragma acc loop reduction(+:Array)
+ for(;;){}
+ }
+
+#pragma acc serial
+ {
+ // expected-error at +2{{OpenACC 'reduction' composite variable must not have non-scalar field}}
+ // expected-note@#COS_FIELD{{invalid field is here}}
+#pragma acc loop reduction(+:ChC)
+ for(;;){}
+ }
+
+#pragma acc serial
+ {
+#pragma acc loop reduction(+:I)
+ for(;;) {
+ // 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(;;) {
+ }
+ }
+ }
+
+#pragma acc serial
+ {
+#pragma acc loop reduction(+:I)
+ for(;;) {
+ // 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(;;) {
+ }
+ }
+ }
+
+#pragma acc serial
+ {
+#pragma acc loop reduction(+:I)
+ for(;;) {
+#pragma acc serial
+ // expected-error at +2{{OpenACC 'reduction' variable must have the same operator in all nested constructs (& vs +)}}
+ // expected-note at -4{{previous clause is here}}
+#pragma acc loop reduction(&:I)
+ for(;;) {
+ }
+ }
+ }
+
+#pragma acc serial reduction(+:I)
+ // expected-error at +2{{OpenACC 'reduction' variable must have the same operator in all nested constructs (& vs +)}}
+ // expected-note at -2{{previous clause is here}}
+#pragma acc loop reduction(&:I)
+ for(;;){}
+
+#pragma acc serial
+#pragma acc loop reduction(&:I)
+ for(;;) {
+ // 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 serial reduction(+:I)
+ ;
+ }
+
+#pragma acc parallel
+ {
+#pragma acc loop reduction(+:I) gang(dim:1)
+ for(;;) {
+ }
+ }
+
+#pragma acc parallel
+ {
+ // expected-error at +2{{OpenACC 'gang' clause with a 'dim' value greater than 1 cannot appear on the same 'loop' construct as a 'reduction' clause}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc loop reduction(+:I) gang(dim:2)
+ for(;;) {
+ }
+ }
+
+#pragma acc parallel
+ {
+ // expected-error at +2{{OpenACC 'reduction' clause cannot appear on the same 'loop' construct as a 'gang' clause with a 'dim' value greater than 1}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc loop gang(dim:2) reduction(+:I)
+ for(;;) {
+ }
+ }
+
+#pragma acc parallel
+ {
+ // expected-error at +2{{OpenACC 'reduction' clause cannot appear on the same 'loop' construct as a 'gang' clause with a 'dim' value greater than 1}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc loop gang gang(dim:1) gang(dim:2) reduction(+:I)
+ for(;;) {
+ }
+ }
+
+#pragma acc parallel num_gangs(1, 2)
+ {
+ // expected-error at +3{{OpenACC 'reduction' clause cannot appear on the same 'loop' construct as a 'gang' clause inside a compute construct with a 'num_gangs' clause with more than one argument}}
+ // expected-note at +2{{previous clause is here}}
+ // expected-note at -4{{previous clause is here}}
+#pragma acc loop gang(dim:1) reduction(+:I)
+ for(;;) {
+ }
+ }
+
+#pragma acc parallel num_gangs(2, 3)
+ {
+ // expected-error at +3{{OpenACC 'gang' clause cannot appear on the same 'loop' construct as a 'reduction' clause inside a compute construct with a 'num_gangs' clause with more than one argument}}
+ // expected-note at +2{{previous clause is here}}
+ // expected-note at -4{{previous clause is here}}
+#pragma acc loop reduction(+:I) gang(dim:1)
+ for(;;) {
+ }
+ }
+}
+
+template<typename IntTy, typename CoSTy, typename ChCTy, unsigned One,
+ unsigned Two>
+void templ_uses() {
+ IntTy I;
+ IntTy Array[5];
+ CoSTy CoS;
+ ChCTy ChC;
+
+#pragma acc serial
+ {
+#pragma acc loop reduction(+:CoS, I)
+ for(;;){}
+ }
+
+#pragma acc serial
+ {
+ // expected-error at +1{{OpenACC 'reduction' variable must be of scalar type, sub-array, or a composite of scalar types; type is 'int[5]'}}
+#pragma acc loop reduction(+:Array)
+ for(;;){}
+ }
+
+#pragma acc serial
+ {
+ // expected-error at +2{{OpenACC 'reduction' composite variable must not have non-scalar field}}
+ // expected-note@#COS_FIELD{{invalid field is here}}
+#pragma acc loop reduction(+:ChC)
+ for(;;){}
+ }
+
+#pragma acc serial
+ {
+#pragma acc loop reduction(+:I)
+ for(;;) {
+ // 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(;;) {
+ }
+ }
+ }
+
+#pragma acc serial
+ {
+#pragma acc loop reduction(+:Array[3])
+ for(;;) {
+ // 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(&:Array[3])
+ for(;;) {
+ }
+ }
+ }
+
+#pragma acc serial
+ {
+#pragma acc loop reduction(+:Array[0:3])
+ for(;;) {
+ // 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(&:Array[1:4])
+ for(;;) {
+ }
+ }
+ }
+
+#pragma acc serial
+ {
+#pragma acc loop reduction(+:I)
+ for(;;) {
+ // 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 serial reduction(&:I)
+ for(;;) {
+ }
+ }
+ }
+
+#pragma acc parallel
+ {
+#pragma acc loop reduction(+:I) gang(dim:One)
+ for(;;) {
+ }
+ }
+
+#pragma acc parallel
+ {
+ // expected-error at +2{{OpenACC 'gang' clause with a 'dim' value greater than 1 cannot appear on the same 'loop' construct as a 'reduction' clause}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc loop reduction(+:I) gang(dim:2)
+ for(;;) {
+ }
+ }
+
+#pragma acc parallel
+ {
+ // expected-error at +2{{OpenACC 'reduction' clause cannot appear on the same 'loop' construct as a 'gang' clause with a 'dim' value greater than 1}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc loop gang(dim:2) reduction(+:I)
+ for(;;) {
+ }
+ }
+#pragma acc parallel
+ {
+ // expected-error at +2{{OpenACC 'gang' clause with a 'dim' value greater than 1 cannot appear on the same 'loop' construct as a 'reduction' clause}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc loop reduction(+:I) gang(dim:Two)
+ for(;;) {
+ }
+ }
+
+#pragma acc parallel
+ {
+ // expected-error at +2{{OpenACC 'reduction' clause cannot appear on the same 'loop' construct as a 'gang' clause with a 'dim' value greater than 1}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc loop gang(dim:Two) reduction(+:I)
+ for(;;) {
+ }
+ }
+
+
+#pragma acc parallel num_gangs(One)
+ {
+#pragma acc loop reduction(+:I) gang(dim:One)
+ for(;;) {
+ }
+ }
+
+#pragma acc parallel num_gangs(Two, 1)
+ {
+ // expected-error at +3{{OpenACC 'gang' clause cannot appear on the same 'loop' construct as a 'reduction' clause inside a compute construct with a 'num_gangs' clause with more than one argument}}
+ // expected-note at +2{{previous clause is here}}
+ // expected-note at -4{{previous clause is here}}
+#pragma acc loop reduction(+:I) gang(dim:One)
+ for(;;) {
+ }
+ }
+
+#pragma acc parallel num_gangs(Two, 1)
+ {
+ // expected-error at +3{{OpenACC 'reduction' clause cannot appear on the same 'loop' construct as a 'gang' clause inside a compute construct with a 'num_gangs' clause with more than one argument}}
+ // expected-note at +2{{previous clause is here}}
+ // expected-note at -4{{previous clause is here}}
+#pragma acc loop gang(dim:One) reduction(+:I)
+ for(;;) {
+ }
+ }
+
+#pragma acc parallel num_gangs(One)
+ {
+#pragma acc loop reduction(+:I) gang(dim:1)
+ for(;;) {
+ }
+ }
+
+#pragma acc parallel num_gangs(Two, 1)
+ {
+ // expected-error at +3{{OpenACC 'gang' clause cannot appear on the same 'loop' construct as a 'reduction' clause inside a compute construct with a 'num_gangs' clause with more than one argument}}
+ // expected-note at +2{{previous clause is here}}
+ // expected-note at -4{{previous clause is here}}
+#pragma acc loop reduction(+:I) gang(dim:1)
+ for(;;) {
+ }
+ }
+
+#pragma acc parallel num_gangs(Two, 1)
+ {
+ // expected-error at +3{{OpenACC 'reduction' clause cannot appear on the same 'loop' construct as a 'gang' clause inside a compute construct with a 'num_gangs' clause with more than one argument}}
+ // expected-note at +2{{previous clause is here}}
+ // expected-note at -4{{previous clause is here}}
+#pragma acc loop gang(dim:1) reduction(+:I)
+ for(;;) {
+ }
+ }
+
+#pragma acc parallel num_gangs(1)
+ {
+#pragma acc loop reduction(+:I) gang(dim:One)
+ for(;;) {
+ }
+ }
+
+#pragma acc parallel num_gangs(2, 1)
+ {
+ // expected-error at +3{{OpenACC 'gang' clause cannot appear on the same 'loop' construct as a 'reduction' clause inside a compute construct with a 'num_gangs' clause with more than one argument}}
+ // expected-note at +2{{previous clause is here}}
+ // expected-note at -4{{previous clause is here}}
+#pragma acc loop reduction(+:I) gang(dim:One)
+ for(;;) {
+ }
+ }
+
+#pragma acc parallel num_gangs(2, 1)
+ {
+ // expected-error at +3{{OpenACC 'reduction' clause cannot appear on the same 'loop' construct as a 'gang' clause inside a compute construct with a 'num_gangs' clause with more than one argument}}
+ // expected-note at +2{{previous clause is here}}
+ // expected-note at -4{{previous clause is here}}
+#pragma acc loop gang(dim:One) reduction(+:I)
+ for(;;) {
+ }
+ }
+}
+
+void inst() {
+ // expected-note at +1{{in instantiation of function template specialization}}
+ templ_uses<int, CompositeOfScalars, CompositeHasComposite, 1, 2>();
+}
+
+
More information about the cfe-commits
mailing list