[clang] 3a4b9f3 - [OpenACC] Implement 'gang' clause for Combined Constructs
via cfe-commits
cfe-commits at lists.llvm.org
Thu Dec 5 06:35:42 PST 2024
Author: erichkeane
Date: 2024-12-05T06:35:36-08:00
New Revision: 3a4b9f38915625c68c78b62de48a3de8b97c5043
URL: https://github.com/llvm/llvm-project/commit/3a4b9f38915625c68c78b62de48a3de8b97c5043
DIFF: https://github.com/llvm/llvm-project/commit/3a4b9f38915625c68c78b62de48a3de8b97c5043.diff
LOG: [OpenACC] Implement 'gang' clause for Combined Constructs
This one is a bit complicated, as it has some interesting interactions,
as 'gang' Sema is required to look at its containing compute construct.
Except in the case of a combined construct, they are the same. This
resulted in a large refactor of the checking code for CheckGangExpr,
plus some additional work on the diagnostics for its interaction with
'num_gangs' and 'vector'/'worker'.
Added:
clang/test/SemaOpenACC/combined-construct-gang-ast.cpp
clang/test/SemaOpenACC/combined-construct-gang-clause.cpp
Modified:
clang/include/clang/AST/OpenACCClause.h
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/combined-construct-device_type-clause.c
clang/test/SemaOpenACC/loop-construct-gang-clause.cpp
Removed:
################################################################################
diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h
index 5ad4c336b6c531..2588c3f645c02b 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -483,6 +483,14 @@ class OpenACCGangClause final
return {getGangKind(I), getExprs()[I]};
}
+ bool hasExprOfKind(OpenACCGangKind GK) const {
+ for (unsigned I = 0; I < getNumExprs(); ++I) {
+ if (getGangKind(I) == GK)
+ return true;
+ }
+ return false;
+ }
+
static OpenACCGangClause *
Create(const ASTContext &Ctx, SourceLocation BeginLoc,
SourceLocation LParenLoc, ArrayRef<OpenACCGangKind> GangKinds,
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 2137cb713164ad..447358f0a5f382 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12755,21 +12755,22 @@ def err_acc_gang_multiple_elt
: Error<"OpenACC 'gang' clause may have at most one %select{unnamed or "
"'num'|'dim'|'static'}0 argument">;
def err_acc_int_arg_invalid
- : Error<"'%1' argument on '%0' clause is not permitted on a%select{n "
- "orphaned|||}2 'loop' construct %select{|associated with a "
- "'parallel' compute construct|associated with a 'kernels' compute "
- "construct|associated with a 'serial' compute construct}2">;
+ : Error<"'%0' argument on '%1' clause is not permitted on a%select{|n "
+ "orphaned}2 '%3' construct%select{| associated with a '%5' compute "
+ "construct}4">;
def err_acc_gang_dim_value
: Error<"argument to 'gang' clause dimension must be %select{a constant "
"expression|1, 2, or 3: evaluated to %1}0">;
def err_acc_num_arg_conflict
- : Error<"'num' argument to '%0' clause not allowed on a 'loop' construct "
- "associated with a 'kernels' construct that has a "
- "'%select{num_gangs|num_workers|vector_length}1' "
- "clause">;
+ : Error<"'num' argument to '%0' clause not allowed on a '%1' "
+ "construct%select{| associated with a '%3' construct}2 that has a "
+ "'%4' clause">;
+def err_acc_num_arg_conflict_reverse
+ : Error<"'num_gangs' clause not allowed on a 'kernels loop' construct that "
+ "has a 'gang' clause with a 'num' argument">;
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">;
+ "clause%select{| on a '%3' 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 "
diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h
index d720cf3c74d87e..a4132534686a81 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -164,9 +164,14 @@ class SemaOpenACC : public SemaBase {
}
/// If there is a current 'active' loop construct with a 'gang' clause on a
- /// 'kernel' construct, this will have the source location for it. This
- /// permits us to implement the restriction of no further 'gang' clauses.
- SourceLocation LoopGangClauseOnKernelLoc;
+ /// 'kernel' construct, this will have the source location for it, and the
+ /// 'kernel kind'. This permits us to implement the restriction of no further
+ /// 'gang' clauses.
+ struct LoopGangOnKernelTy {
+ SourceLocation Loc;
+ OpenACCDirectiveKind DirKind = OpenACCDirectiveKind::Invalid;
+ } LoopGangClauseOnKernel;
+
/// If there is a current 'active' loop construct with a 'worker' clause on it
/// (on any sort of construct), this has the source location for it. This
/// permits us to implement the restriction of no further 'gang' or 'worker'
@@ -705,7 +710,9 @@ class SemaOpenACC : public SemaBase {
ExprResult CheckTileSizeExpr(Expr *SizeExpr);
// Check a single expression on a gang clause.
- ExprResult CheckGangExpr(OpenACCGangKind GK, Expr *E);
+ ExprResult CheckGangExpr(ArrayRef<const OpenACCClause *> ExistingClauses,
+ OpenACCDirectiveKind DK, OpenACCGangKind GK,
+ Expr *E);
// Does the checking for a 'gang' clause that needs to be done in dependent
// and not dependent cases.
@@ -771,7 +778,7 @@ class SemaOpenACC : public SemaBase {
SemaOpenACC &SemaRef;
ComputeConstructInfo OldActiveComputeConstructInfo;
OpenACCDirectiveKind DirKind;
- SourceLocation OldLoopGangClauseOnKernelLoc;
+ LoopGangOnKernelTy OldLoopGangClauseOnKernel;
SourceLocation OldLoopWorkerClauseLoc;
SourceLocation OldLoopVectorClauseLoc;
LoopWithoutSeqCheckingInfo OldLoopWithoutSeqInfo;
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 30491d6f3a3bd6..16348d0b8837ef 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -737,6 +737,24 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitNumGangsClause(
return nullptr;
}
}
+
+ // OpenACC 3.3 Section 2.9.2:
+ // An argument with no keyword or with the 'num' wkeyword 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(
+ ExistingClauses, llvm::IsaPred<OpenACCGangClause>);
+
+ for (auto *GC : GangClauses) {
+ if (cast<OpenACCGangClause>(GC)->hasExprOfKind(OpenACCGangKind::Num)) {
+ SemaRef.Diag(Clause.getBeginLoc(),
+ diag::err_acc_num_arg_conflict_reverse);
+ SemaRef.Diag(GC->getBeginLoc(), diag::note_acc_previous_clause_here);
+ return nullptr;
+ }
+ }
+ }
+
return OpenACCNumGangsClause::Create(
Ctx, Clause.getBeginLoc(), Clause.getLParenLoc(), Clause.getIntExprs(),
Clause.getEndLoc());
@@ -1033,6 +1051,136 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitIndependentClause(
Clause.getEndLoc());
}
+ExprResult CheckGangStaticExpr(SemaOpenACC &S, Expr *E) {
+ if (isa<OpenACCAsteriskSizeExpr>(E))
+ return E;
+ return S.ActOnIntExpr(OpenACCDirectiveKind::Invalid, OpenACCClauseKind::Gang,
+ E->getBeginLoc(), E);
+}
+
+bool IsOrphanLoop(OpenACCDirectiveKind DK, OpenACCDirectiveKind AssocKind) {
+ return DK == OpenACCDirectiveKind::Loop &&
+ AssocKind == OpenACCDirectiveKind::Invalid;
+}
+
+bool HasAssocKind(OpenACCDirectiveKind DK, OpenACCDirectiveKind AssocKind) {
+ return DK == OpenACCDirectiveKind::Loop &&
+ AssocKind != OpenACCDirectiveKind::Invalid;
+}
+
+ExprResult DiagIntArgInvalid(SemaOpenACC &S, Expr *E, OpenACCGangKind GK,
+ OpenACCClauseKind CK, OpenACCDirectiveKind DK,
+ OpenACCDirectiveKind AssocKind) {
+ S.Diag(E->getBeginLoc(), diag::err_acc_int_arg_invalid)
+ << GK << CK << IsOrphanLoop(DK, AssocKind) << DK
+ << HasAssocKind(DK, AssocKind) << AssocKind;
+ return ExprError();
+}
+
+ExprResult CheckGangParallelExpr(SemaOpenACC &S, OpenACCDirectiveKind DK,
+ OpenACCDirectiveKind AssocKind,
+ OpenACCGangKind GK, Expr *E) {
+ switch (GK) {
+ case OpenACCGangKind::Static:
+ return CheckGangStaticExpr(S, E);
+ case OpenACCGangKind::Num:
+ // OpenACC 3.3 2.9.2: When the parent compute construct is a parallel
+ // construct, or an orphaned loop construct, the gang clause behaves as
+ // follows. ... The num argument is not allowed.
+ return DiagIntArgInvalid(S, E, GK, OpenACCClauseKind::Gang, DK, AssocKind);
+ case OpenACCGangKind::Dim: {
+ // OpenACC 3.3 2.9.2: When the parent compute construct is a parallel
+ // construct, or an orphaned loop construct, the gang clause behaves as
+ // follows. ... The dim argument must be a constant positive integer value
+ // 1, 2, or 3.
+ if (!E)
+ return ExprError();
+ ExprResult Res =
+ S.ActOnIntExpr(OpenACCDirectiveKind::Invalid, OpenACCClauseKind::Gang,
+ E->getBeginLoc(), E);
+
+ if (!Res.isUsable())
+ return Res;
+
+ if (Res.get()->isInstantiationDependent())
+ return Res;
+
+ std::optional<llvm::APSInt> ICE =
+ Res.get()->getIntegerConstantExpr(S.getASTContext());
+
+ if (!ICE || *ICE <= 0 || ICE > 3) {
+ S.Diag(Res.get()->getBeginLoc(), diag::err_acc_gang_dim_value)
+ << ICE.has_value() << ICE.value_or(llvm::APSInt{}).getExtValue();
+ return ExprError();
+ }
+
+ return ExprResult{
+ ConstantExpr::Create(S.getASTContext(), Res.get(), APValue{*ICE})};
+ }
+ }
+ llvm_unreachable("Unknown gang kind in gang parallel check");
+}
+
+ExprResult CheckGangKernelsExpr(SemaOpenACC &S,
+ ArrayRef<const OpenACCClause *> ExistingClauses,
+ OpenACCDirectiveKind DK,
+ OpenACCDirectiveKind AssocKind,
+ OpenACCGangKind GK, Expr *E) {
+ switch (GK) {
+ // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels
+ // construct, the gang clause behaves as follows. ... The dim argument is
+ // not allowed.
+ case OpenACCGangKind::Dim:
+ return DiagIntArgInvalid(S, E, GK, OpenACCClauseKind::Gang, DK, AssocKind);
+ case OpenACCGangKind::Num: {
+ // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels
+ // construct, the gang clause behaves as follows. ... An argument with no
+ // keyword or with num keyword is only allowed when num_gangs does not
+ // appear on the kernels construct. ... The region of a loop with the gang
+ // clause may not contain another loop with a gang clause unless within a
+ // nested compute region.
+
+ // If this is a 'combined' construct, search the list of existing clauses.
+ // Else we need to search the containing 'kernel'.
+ auto Collection = isOpenACCCombinedDirectiveKind(DK)
+ ? ExistingClauses
+ : S.getActiveComputeConstructInfo().Clauses;
+
+ const auto *Itr =
+ llvm::find_if(Collection, llvm::IsaPred<OpenACCNumGangsClause>);
+
+ if (Itr != Collection.end()) {
+ S.Diag(E->getBeginLoc(), diag::err_acc_num_arg_conflict)
+ << OpenACCClauseKind::Gang << DK << HasAssocKind(DK, AssocKind)
+ << AssocKind << OpenACCClauseKind::NumGangs;
+
+ S.Diag((*Itr)->getBeginLoc(), diag::note_acc_previous_clause_here);
+ return ExprError();
+ }
+ return ExprResult{E};
+ }
+ case OpenACCGangKind::Static:
+ return CheckGangStaticExpr(S, E);
+ return ExprError();
+ }
+ llvm_unreachable("Unknown gang kind in gang kernels check");
+}
+
+ExprResult CheckGangSerialExpr(SemaOpenACC &S, OpenACCDirectiveKind DK,
+ OpenACCDirectiveKind AssocKind,
+ OpenACCGangKind GK, Expr *E) {
+ switch (GK) {
+ // 'dim' and 'num' don't really make sense on serial, and GCC rejects them
+ // too, so we disallow them too.
+ case OpenACCGangKind::Dim:
+ case OpenACCGangKind::Num:
+ return DiagIntArgInvalid(S, E, GK, OpenACCClauseKind::Gang, DK, AssocKind);
+ case OpenACCGangKind::Static:
+ return CheckGangStaticExpr(S, E);
+ }
+ llvm_unreachable("Unknown gang kind in gang serial check");
+}
+
OpenACCClause *SemaOpenACCClauseVisitor::VisitVectorClause(
SemaOpenACC::OpenACCParsedClause &Clause) {
if (DiagIfSeqClause(Clause))
@@ -1054,8 +1202,9 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitVectorClause(
case OpenACCDirectiveKind::Serial:
// GCC disallows this, and there is no real good reason for us to permit
// it, so disallow until we come up with a use case that makes sense.
- SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_int_arg_invalid)
- << OpenACCClauseKind::Vector << "num" << /*serial=*/3;
+ DiagIntArgInvalid(SemaRef, IntExpr, OpenACCGangKind::Num,
+ OpenACCClauseKind::Vector, Clause.getDirectiveKind(),
+ SemaRef.getActiveComputeConstructInfo().Kind);
IntExpr = nullptr;
break;
case OpenACCDirectiveKind::Kernels: {
@@ -1064,7 +1213,11 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitVectorClause(
llvm::IsaPred<OpenACCVectorLengthClause>);
if (Itr != SemaRef.getActiveComputeConstructInfo().Clauses.end()) {
SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict)
- << OpenACCClauseKind::Vector << /*vector_length=*/2;
+ << OpenACCClauseKind::Vector << Clause.getDirectiveKind()
+ << HasAssocKind(Clause.getDirectiveKind(),
+ SemaRef.getActiveComputeConstructInfo().Kind)
+ << SemaRef.getActiveComputeConstructInfo().Kind
+ << OpenACCClauseKind::VectorLength;
SemaRef.Diag((*Itr)->getBeginLoc(),
diag::note_acc_previous_clause_here);
@@ -1113,18 +1266,11 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitWorkerClause(
if (IntExpr) {
switch (SemaRef.getActiveComputeConstructInfo().Kind) {
case OpenACCDirectiveKind::Invalid:
- SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_int_arg_invalid)
- << OpenACCClauseKind::Worker << "num" << /*orphan=*/0;
- IntExpr = nullptr;
- break;
case OpenACCDirectiveKind::Parallel:
- SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_int_arg_invalid)
- << OpenACCClauseKind::Worker << "num" << /*parallel=*/1;
- IntExpr = nullptr;
- break;
case OpenACCDirectiveKind::Serial:
- SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_int_arg_invalid)
- << OpenACCClauseKind::Worker << "num" << /*serial=*/3;
+ DiagIntArgInvalid(SemaRef, IntExpr, OpenACCGangKind::Num,
+ OpenACCClauseKind::Worker, Clause.getDirectiveKind(),
+ SemaRef.getActiveComputeConstructInfo().Kind);
IntExpr = nullptr;
break;
case OpenACCDirectiveKind::Kernels: {
@@ -1133,7 +1279,11 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitWorkerClause(
llvm::IsaPred<OpenACCNumWorkersClause>);
if (Itr != SemaRef.getActiveComputeConstructInfo().Clauses.end()) {
SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict)
- << OpenACCClauseKind::Worker << /*num_workers=*/1;
+ << OpenACCClauseKind::Worker << Clause.getDirectiveKind()
+ << HasAssocKind(Clause.getDirectiveKind(),
+ SemaRef.getActiveComputeConstructInfo().Kind)
+ << SemaRef.getActiveComputeConstructInfo().Kind
+ << OpenACCClauseKind::NumWorkers;
SemaRef.Diag((*Itr)->getBeginLoc(),
diag::note_acc_previous_clause_here);
@@ -1187,12 +1337,15 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitGangClause(
// Restrictions only properly implemented on 'loop' constructs, and it is
// the only construct that can do anything with this, so skip/treat as
// unimplemented for the combined constructs.
- if (Clause.getDirectiveKind() != OpenACCDirectiveKind::Loop)
+ if (Clause.getDirectiveKind() != OpenACCDirectiveKind::Loop &&
+ !isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind()))
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.
+ // 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) {
@@ -1229,31 +1382,13 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitGangClause(
for (unsigned I = 0; I < Clause.getIntExprs().size(); ++I) {
OpenACCGangKind GK = Clause.getGangKinds()[I];
- ExprResult ER = SemaRef.CheckGangExpr(GK, Clause.getIntExprs()[I]);
+ ExprResult ER =
+ SemaRef.CheckGangExpr(ExistingClauses, Clause.getDirectiveKind(), GK,
+ Clause.getIntExprs()[I]);
if (!ER.isUsable())
continue;
- // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels
- // construct, the gang clause behaves as follows. ... An argument with no
- // keyword or with num keyword is only allowed when num_gangs does not
- // appear on the kernels construct.
- if (SemaRef.getActiveComputeConstructInfo().Kind ==
- OpenACCDirectiveKind::Kernels &&
- GK == OpenACCGangKind::Num) {
- const auto *Itr =
- llvm::find_if(SemaRef.getActiveComputeConstructInfo().Clauses,
- llvm::IsaPred<OpenACCNumGangsClause>);
-
- if (Itr != SemaRef.getActiveComputeConstructInfo().Clauses.end()) {
- SemaRef.Diag(ER.get()->getBeginLoc(), diag::err_acc_num_arg_conflict)
- << OpenACCClauseKind::Gang << /*num_gangs=*/0;
- SemaRef.Diag((*Itr)->getBeginLoc(),
- diag::note_acc_previous_clause_here);
- continue;
- }
- }
-
// OpenACC 3.3 2.9: 'gang-arg-list' may have at most one num, one dim, and
// one static argument.
if (ExistingElemLoc[static_cast<unsigned>(GK)].isValid()) {
@@ -1269,47 +1404,50 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitGangClause(
IntExprs.push_back(ER.get());
}
- // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels
- // construct, the gang clause behaves as follows. ... The region of a loop
- // with a gang clause may not contain another loop with a gang clause unless
- // within a nested compute region.
- if (SemaRef.LoopGangClauseOnKernelLoc.isValid()) {
- // This handles the 'inner loop' diagnostic, but we cannot set that we're on
- // one of these until we get to the end of the construct.
- SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_in_clause_region)
- << OpenACCClauseKind::Gang << OpenACCClauseKind::Gang
- << /*kernels construct info*/ 1;
- SemaRef.Diag(SemaRef.LoopGangClauseOnKernelLoc,
- diag::note_acc_previous_clause_here);
- return nullptr;
- }
+ if (!isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind())) {
+ // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels
+ // construct, the gang clause behaves as follows. ... The region of a loop
+ // with a gang clause may not contain another loop with a gang clause unless
+ // within a nested compute region.
+ if (SemaRef.LoopGangClauseOnKernel.Loc.isValid()) {
+ // This handles the 'inner loop' diagnostic, but we cannot set that we're
+ // on one of these until we get to the end of the construct.
+ SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_in_clause_region)
+ << OpenACCClauseKind::Gang << OpenACCClauseKind::Gang
+ << /*kernels construct info*/ 1
+ << SemaRef.LoopGangClauseOnKernel.DirKind;
+ SemaRef.Diag(SemaRef.LoopGangClauseOnKernel.Loc,
+ diag::note_acc_previous_clause_here);
+ return nullptr;
+ }
- // OpenACC 3.3 2.9.3: The region of a loop with a 'worker' clause may not
- // contain a loop with a gang or worker clause unless within a nested compute
- // region.
- if (SemaRef.LoopWorkerClauseLoc.isValid()) {
- // This handles the 'inner loop' diagnostic, but we cannot set that we're on
- // one of these until we get to the end of the construct.
- SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_in_clause_region)
- << OpenACCClauseKind::Gang << OpenACCClauseKind::Worker
- << /*kernels construct info*/ 1;
- SemaRef.Diag(SemaRef.LoopWorkerClauseLoc,
- diag::note_acc_previous_clause_here);
- return nullptr;
- }
+ // OpenACC 3.3 2.9.3: The region of a loop with a 'worker' clause may not
+ // contain a loop with a gang or worker clause unless within a nested
+ // compute region.
+ if (SemaRef.LoopWorkerClauseLoc.isValid()) {
+ // This handles the 'inner loop' diagnostic, but we cannot set that we're
+ // on one of these until we get to the end of the construct.
+ SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_in_clause_region)
+ << OpenACCClauseKind::Gang << OpenACCClauseKind::Worker
+ << /*!kernels construct info*/ 0;
+ SemaRef.Diag(SemaRef.LoopWorkerClauseLoc,
+ diag::note_acc_previous_clause_here);
+ return nullptr;
+ }
- // OpenACC 3.3 2.9.4: The region of a loop with a 'vector' clause may not
- // contain a loop with a gang, worker, or vector clause unless within a nested
- // compute region.
- if (SemaRef.LoopVectorClauseLoc.isValid()) {
- // This handles the 'inner loop' diagnostic, but we cannot set that we're on
- // one of these until we get to the end of the construct.
- SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_in_clause_region)
- << OpenACCClauseKind::Gang << OpenACCClauseKind::Vector
- << /*kernels construct info*/ 1;
- SemaRef.Diag(SemaRef.LoopVectorClauseLoc,
- diag::note_acc_previous_clause_here);
- return nullptr;
+ // OpenACC 3.3 2.9.4: The region of a loop with a 'vector' clause may not
+ // contain a loop with a gang, worker, or vector clause unless within a
+ // nested compute region.
+ if (SemaRef.LoopVectorClauseLoc.isValid()) {
+ // This handles the 'inner loop' diagnostic, but we cannot set that we're
+ // on one of these until we get to the end of the construct.
+ SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_in_clause_region)
+ << OpenACCClauseKind::Gang << OpenACCClauseKind::Vector
+ << /*!kernels construct info*/ 0;
+ SemaRef.Diag(SemaRef.LoopVectorClauseLoc,
+ diag::note_acc_previous_clause_here);
+ return nullptr;
+ }
}
return SemaRef.CheckGangClause(ExistingClauses, Clause.getBeginLoc(),
@@ -1489,7 +1627,7 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
ArrayRef<const OpenACCClause *> UnInstClauses,
ArrayRef<OpenACCClause *> Clauses)
: SemaRef(S), OldActiveComputeConstructInfo(S.ActiveComputeConstructInfo),
- DirKind(DK), OldLoopGangClauseOnKernelLoc(S.LoopGangClauseOnKernelLoc),
+ DirKind(DK), OldLoopGangClauseOnKernel(S.LoopGangClauseOnKernel),
OldLoopWorkerClauseLoc(S.LoopWorkerClauseLoc),
OldLoopVectorClauseLoc(S.LoopVectorClauseLoc),
OldLoopWithoutSeqInfo(S.LoopWithoutSeqInfo),
@@ -1510,7 +1648,7 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
// within a nested compute region.
//
// Implement the 'unless within a nested compute region' part.
- SemaRef.LoopGangClauseOnKernelLoc = {};
+ SemaRef.LoopGangClauseOnKernel = {};
SemaRef.LoopWorkerClauseLoc = {};
SemaRef.LoopVectorClauseLoc = {};
SemaRef.LoopWithoutSeqInfo = {};
@@ -1524,8 +1662,7 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
SetCollapseInfoBeforeAssociatedStmt(UnInstClauses, Clauses);
SetTileInfoBeforeAssociatedStmt(UnInstClauses, Clauses);
- // TODO: OpenACC: We need to set these 3, CollapseInfo, and TileInfo
- SemaRef.LoopGangClauseOnKernelLoc = {};
+ SemaRef.LoopGangClauseOnKernel = {};
SemaRef.LoopWorkerClauseLoc = {};
SemaRef.LoopVectorClauseLoc = {};
@@ -1548,7 +1685,7 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
// This handles the 'outer loop' part of this.
auto *Itr = llvm::find_if(Clauses, llvm::IsaPred<OpenACCGangClause>);
if (Itr != Clauses.end())
- SemaRef.LoopGangClauseOnKernelLoc = (*Itr)->getBeginLoc();
+ SemaRef.LoopGangClauseOnKernel = {(*Itr)->getBeginLoc(), DirKind};
}
if (UnInstClauses.empty()) {
@@ -1586,7 +1723,8 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
// This handles the 'outer loop' part of this.
auto *Itr = llvm::find_if(Clauses, llvm::IsaPred<OpenACCGangClause>);
if (Itr != Clauses.end())
- SemaRef.LoopGangClauseOnKernelLoc = (*Itr)->getBeginLoc();
+ SemaRef.LoopGangClauseOnKernel = {(*Itr)->getBeginLoc(),
+ OpenACCDirectiveKind::Kernels};
}
if (UnInstClauses.empty()) {
@@ -1671,7 +1809,7 @@ void SemaOpenACC::AssociatedStmtRAII::SetTileInfoBeforeAssociatedStmt(
SemaOpenACC::AssociatedStmtRAII::~AssociatedStmtRAII() {
SemaRef.ActiveComputeConstructInfo = OldActiveComputeConstructInfo;
- SemaRef.LoopGangClauseOnKernelLoc = OldLoopGangClauseOnKernelLoc;
+ SemaRef.LoopGangClauseOnKernel = OldLoopGangClauseOnKernel;
SemaRef.LoopWorkerClauseLoc = OldLoopWorkerClauseLoc;
SemaRef.LoopVectorClauseLoc = OldLoopVectorClauseLoc;
SemaRef.LoopWithoutSeqInfo = OldLoopWithoutSeqInfo;
@@ -2328,106 +2466,47 @@ ExprResult SemaOpenACC::CheckCollapseLoopCount(Expr *LoopCount) {
ConstantExpr::Create(getASTContext(), LoopCount, APValue{*ICE})};
}
-namespace {
-ExprResult CheckGangStaticExpr(SemaOpenACC &S, Expr *E) {
- if (isa<OpenACCAsteriskSizeExpr>(E))
- return E;
- return S.ActOnIntExpr(OpenACCDirectiveKind::Invalid, OpenACCClauseKind::Gang,
- E->getBeginLoc(), E);
-}
-} // namespace
-
-ExprResult SemaOpenACC::CheckGangExpr(OpenACCGangKind GK, Expr *E) {
- // Gang Expr legality depends on the associated compute construct.
- switch (ActiveComputeConstructInfo.Kind) {
- case OpenACCDirectiveKind::Invalid:
- case OpenACCDirectiveKind::Parallel: {
- switch (GK) {
- // OpenACC 3.3 2.9.2: When the parent compute construct is a parallel
- // construct, or an orphaned loop construct, the gang clause behaves as
- // follows. ... The dim argument must be a constant positive integer value
- // 1, 2, or 3.
- case OpenACCGangKind::Dim: {
- if (!E)
- return ExprError();
- ExprResult Res =
- ActOnIntExpr(OpenACCDirectiveKind::Invalid, OpenACCClauseKind::Gang,
- E->getBeginLoc(), E);
-
- if (!Res.isUsable())
- return Res;
-
- if (Res.get()->isInstantiationDependent())
- return Res;
-
- std::optional<llvm::APSInt> ICE =
- Res.get()->getIntegerConstantExpr(getASTContext());
-
- if (!ICE || *ICE <= 0 || ICE > 3) {
- Diag(Res.get()->getBeginLoc(), diag::err_acc_gang_dim_value)
- << ICE.has_value() << ICE.value_or(llvm::APSInt{}).getExtValue();
- return ExprError();
- }
-
- return ExprResult{
- ConstantExpr::Create(getASTContext(), Res.get(), APValue{*ICE})};
- }
- // OpenACC 3.3 2.9.2: When the parent compute construct is a parallel
- // construct, or an orphaned loop construct, the gang clause behaves as
- // follows. ... The num argument is not allowed.
- case OpenACCGangKind::Num:
- Diag(E->getBeginLoc(), diag::err_acc_int_arg_invalid)
- << OpenACCClauseKind::Gang << GK
- << (/*orphan/parallel=*/ActiveComputeConstructInfo.Kind ==
- OpenACCDirectiveKind::Parallel
- ? 1
- : 0);
- return ExprError();
- case OpenACCGangKind::Static:
- return CheckGangStaticExpr(*this, E);
- }
- } break;
- case OpenACCDirectiveKind::Kernels: {
- switch (GK) {
- // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels
- // construct, the gang clause behaves as follows. ... The dim argument is
- // not allowed.
- case OpenACCGangKind::Dim:
- Diag(E->getBeginLoc(), diag::err_acc_int_arg_invalid)
- << OpenACCClauseKind::Gang << GK << /*kernels=*/2;
- return ExprError();
- // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels
- // construct, the gang clause behaves as follows. ... An argument with no
- // keyword or with num keyword is only allowed when num_gangs does not
- // appear on the kernels construct. ... The region of a loop with the gang
- // clause may not contain another loop with a gang clause unless within a
- // nested compute region.
- case OpenACCGangKind::Num:
- // This isn't allowed if there is a 'num_gangs' on the kernel construct,
- // and makes loop-with-gang-clause ill-formed inside of this 'loop', but
- // nothing can be enforced here.
- return ExprResult{E};
- case OpenACCGangKind::Static:
- return CheckGangStaticExpr(*this, E);
- }
- } break;
- case OpenACCDirectiveKind::Serial: {
- switch (GK) {
- // 'dim' and 'num' don't really make sense on serial, and GCC rejects them
- // too, so we disallow them too.
- case OpenACCGangKind::Dim:
- case OpenACCGangKind::Num:
- Diag(E->getBeginLoc(), diag::err_acc_int_arg_invalid)
- << OpenACCClauseKind::Gang << GK << /*Kernels=*/3;
- return ExprError();
- case OpenACCGangKind::Static:
- return CheckGangStaticExpr(*this, E);
+ExprResult
+SemaOpenACC::CheckGangExpr(ArrayRef<const OpenACCClause *> ExistingClauses,
+ OpenACCDirectiveKind DK, OpenACCGangKind GK,
+ Expr *E) {
+ // There are two cases for the enforcement here: the 'current' directive is a
+ // 'loop', where we need to check the active compute construct kind, or the
+ // current directive is a 'combined' construct, where we have to check the
+ // current one.
+ switch (DK) {
+ case OpenACCDirectiveKind::ParallelLoop:
+ return CheckGangParallelExpr(*this, DK, ActiveComputeConstructInfo.Kind, GK,
+ E);
+ case OpenACCDirectiveKind::SerialLoop:
+ return CheckGangSerialExpr(*this, DK, ActiveComputeConstructInfo.Kind, GK,
+ E);
+ case OpenACCDirectiveKind::KernelsLoop:
+ return CheckGangKernelsExpr(*this, ExistingClauses, DK,
+ ActiveComputeConstructInfo.Kind, GK, E);
+ case OpenACCDirectiveKind::Loop:
+ switch (ActiveComputeConstructInfo.Kind) {
+ case OpenACCDirectiveKind::Invalid:
+ case OpenACCDirectiveKind::Parallel:
+ case OpenACCDirectiveKind::ParallelLoop:
+ return CheckGangParallelExpr(*this, DK, ActiveComputeConstructInfo.Kind,
+ GK, E);
+ case OpenACCDirectiveKind::SerialLoop:
+ case OpenACCDirectiveKind::Serial:
+ return CheckGangSerialExpr(*this, DK, ActiveComputeConstructInfo.Kind, GK,
+ E);
+ case OpenACCDirectiveKind::KernelsLoop:
+ case OpenACCDirectiveKind::Kernels:
+ return CheckGangKernelsExpr(*this, ExistingClauses, DK,
+ ActiveComputeConstructInfo.Kind, GK, E);
+ default:
+ llvm_unreachable("Non compute construct in active compute construct?");
}
- } break;
default:
- llvm_unreachable("Non compute construct in active compute construct?");
+ // TODO: OpenACC: when we implement this on 'routine', we'll have to
+ // implement its checking here.
+ llvm_unreachable("Invalid directive kind for a Gang clause");
}
-
llvm_unreachable("Compute construct directive not handled?");
}
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index 3a8f2d95f329b8..81e515e7cb2a9a 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -12027,7 +12027,9 @@ void OpenACCClauseTransform<Derived>::VisitGangClause(
if (!ER.isUsable())
continue;
- ER = Self.getSema().OpenACC().CheckGangExpr(C.getExpr(I).first, ER.get());
+ ER = Self.getSema().OpenACC().CheckGangExpr(ExistingClauses,
+ ParsedClause.getDirectiveKind(),
+ C.getExpr(I).first, ER.get());
if (!ER.isUsable())
continue;
TransformedGangKinds.push_back(C.getExpr(I).first);
diff --git a/clang/test/AST/ast-print-openacc-combined-construct.cpp b/clang/test/AST/ast-print-openacc-combined-construct.cpp
index 435c770c7457d1..40f174e539c01e 100644
--- a/clang/test/AST/ast-print-openacc-combined-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-combined-construct.cpp
@@ -242,4 +242,75 @@ void foo() {
#pragma acc parallel loop vector_length((int)array[1])
for(int i = 0;i<5;++i);
+// CHECK: #pragma acc parallel loop gang(dim: 2)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc parallel loop gang(dim:2)
+ for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc serial loop gang(static: i)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc serial loop gang(static:i)
+ for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc parallel loop gang(static: i) gang(dim: 2)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc parallel loop gang(static:i) gang(dim:2)
+ for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc parallel loop gang(static: i, dim: 2)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc parallel loop gang(static:i, dim:2)
+ for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc parallel loop gang(dim: 2)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc parallel loop gang(dim:2)
+ for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc parallel loop gang(static: i)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc parallel loop gang(static:i)
+ for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc parallel loop gang(static: i) gang(dim: 2)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc parallel loop gang(static:i) gang(dim:2)
+ for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc parallel loop gang(static: i, dim: 2)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc parallel loop gang(static:i, dim:2)
+ for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc kernels loop gang(num: i) gang(static: i)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc kernels loop gang(i) gang(static:i)
+ for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc kernels loop gang(num: i) gang(static: i)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc kernels loop gang(num:i) gang(static:i)
+ for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc serial loop gang(static: i)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc serial loop gang(static:i)
+ for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc serial loop gang(static: *)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc serial loop gang(static:*)
+ 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 a6f57a63a91ddf..e35bd6da2f18b3 100644
--- a/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
+++ b/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
@@ -155,7 +155,6 @@ void uses() {
#pragma acc parallel loop auto tile(1+2, 1)
for(unsigned j = 0; j < 5; ++j)
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'gang' not yet implemented}}
#pragma acc parallel loop auto gang
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop auto wait
@@ -276,7 +275,6 @@ void uses() {
#pragma acc parallel loop tile(1+2, 1) auto
for(unsigned j = 0; j < 5; ++j)
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'gang' not yet implemented}}
#pragma acc parallel loop gang auto
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop wait auto
@@ -398,7 +396,6 @@ void uses() {
#pragma acc parallel loop independent tile(1+2, 1)
for(unsigned j = 0; j < 5; ++j)
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'gang' not yet implemented}}
#pragma acc parallel loop independent gang
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop independent wait
@@ -519,7 +516,6 @@ void uses() {
#pragma acc parallel loop tile(1+2, 1) independent
for(unsigned j = 0; j < 5; ++j)
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'gang' not yet implemented}}
#pragma acc parallel loop gang independent
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop wait independent
@@ -650,9 +646,8 @@ void uses() {
#pragma acc parallel loop seq wait
for(unsigned i = 0; i < 5; ++i);
- // TODOexpected-error at +2{{OpenACC clause 'seq' may not appear on the same construct as a 'gang' clause on a 'parallel loop' construct}}
- // TODOexpected-note at +1{{previous clause is here}}
- // expected-warning at +1{{OpenACC clause 'gang' not yet implemented}}
+ // expected-error at +2{{OpenACC clause 'seq' may not appear on the same construct as a 'gang' clause on a 'parallel loop' construct}}
+ // expected-note at +1{{previous clause is here}}
#pragma acc parallel loop gang seq
for(unsigned i = 0; i < 5; ++i);
// TODOexpected-error at +2{{OpenACC clause 'seq' may not appear on the same construct as a 'worker' clause on a 'parallel loop' construct}}
diff --git a/clang/test/SemaOpenACC/combined-construct-device_type-clause.c b/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
index 9a60fb4c665e5a..d3ed8234d16b14 100644
--- a/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
@@ -212,7 +212,6 @@ void uses() {
for(int j = 0; j < 5; ++j)
for(int i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
#pragma acc serial loop dtype(*) gang
for(int i = 0; i < 5; ++i);
#pragma acc parallel loop device_type(*) wait
diff --git a/clang/test/SemaOpenACC/combined-construct-gang-ast.cpp b/clang/test/SemaOpenACC/combined-construct-gang-ast.cpp
new file mode 100644
index 00000000000000..f179b928215e71
--- /dev/null
+++ b/clang/test/SemaOpenACC/combined-construct-gang-ast.cpp
@@ -0,0 +1,216 @@
+
+// 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 NormalUses() {
+ // CHECK: FunctionDecl{{.*}}NormalUses
+ // CHECK-NEXT: CompoundStmt
+
+ int Val;
+
+#pragma acc parallel loop gang(dim:1)
+ for(int i = 0; i < 5; ++i);
+ // CHECK: OpenACCCombinedConstruct{{.*}} parallel loop
+ // CHECK-NEXT: gang clause dim
+ // CHECK-NEXT: ConstantExpr
+ // CHECK-NEXT: value: Int 1
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+
+#pragma acc parallel loop gang(static:Val)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+ // CHECK-NEXT: gang clause static
+ // CHECK-NEXT: ImplicitCastExpr
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+
+#pragma acc kernels loop gang(num:1) gang(static:Val)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
+ // CHECK-NEXT: gang clause num
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1
+ // CHECK-NEXT: gang clause static
+ // CHECK-NEXT: ImplicitCastExpr
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+
+#pragma acc parallel loop gang(dim:1, static:Val)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+ // CHECK-NEXT: gang clause dim static
+ // CHECK-NEXT: ConstantExpr
+ // CHECK-NEXT: value: Int 1
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1
+ // CHECK-NEXT: ImplicitCastExpr
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+
+#pragma acc serial loop gang(static:Val)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} serial loop
+ // CHECK-NEXT: gang clause static
+ // CHECK-NEXT: ImplicitCastExpr
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+
+#pragma acc serial loop gang(static:*)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} serial loop
+ // CHECK-NEXT: gang clause static
+ // CHECK-NEXT: OpenACCAsteriskSizeExpr
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+
+#pragma acc serial loop gang
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} serial loop
+ // CHECK-NEXT: gang clause
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+
+#pragma acc kernels loop gang(num:1)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
+ // CHECK-NEXT: gang clause num
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+}
+
+template<typename T, unsigned One>
+void TemplateUses(T Val) {
+ // CHECK: FunctionTemplateDecl{{.*}}TemplateUses
+ // CHECK-NEXT: TemplateTypeParmDecl {{.*}} referenced typename depth 0 index 0 T
+ // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} referenced 'unsigned int' depth 0 index 1 One
+ // CHECK-NEXT: FunctionDecl{{.*}} TemplateUses 'void (T)'
+ // CHECK-NEXT: ParmVarDecl{{.*}} referenced Val 'T'
+ // CHECK-NEXT: CompoundStmt
+
+#pragma acc parallel loop gang(dim:One)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+ // CHECK-NEXT: gang clause dim
+ // CHECK-NEXT: DeclRefExpr{{.*}}'One' 'unsigned int'
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+
+#pragma acc parallel loop gang(static:Val)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+ // CHECK-NEXT: gang clause static
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'T'
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+
+
+#pragma acc parallel loop gang(static:*)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+ // CHECK-NEXT: gang clause static
+ // CHECK-NEXT: OpenACCAsteriskSizeExpr
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+
+#pragma acc parallel loop gang(dim:One) gang(static:Val)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+ // CHECK-NEXT: gang clause dim
+ // CHECK-NEXT: DeclRefExpr{{.*}}'One' 'unsigned int'
+ // CHECK-NEXT: gang clause static
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'T'
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+
+#pragma acc parallel loop gang(dim:One, static:Val)
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+ // CHECK-NEXT: gang clause dim static
+ // CHECK-NEXT: DeclRefExpr{{.*}}'One' 'unsigned int'
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'T'
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+
+#pragma acc serial loop gang
+ for(int i = 0; i < 5; ++i);
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} serial loop
+ // CHECK-NEXT: gang clause
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+
+ // Instantiation:
+ // CHECK-NEXT: FunctionDecl{{.*}} used TemplateUses 'void (int)' implicit_instantiation
+ // CHECK-NEXT: TemplateArgument type 'int'
+ // CHECK-NEXT: BuiltinType{{.*}} 'int'
+ // CHECK-NEXT: TemplateArgument integral '1U'
+ // CHECK-NEXT: ParmVarDecl{{.*}} used Val 'int'
+ // CHECK-NEXT: CompoundStmt
+
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+ // CHECK-NEXT: gang clause dim
+ // CHECK-NEXT: ConstantExpr
+ // CHECK-NEXT: value: Int 1
+ // CHECK-NEXT: SubstNonTypeTemplateParmExpr
+ // CHECK-NEXT: NonTypeTemplateParmDecl
+ // CHECK-NEXT: IntegerLiteral{{.*}}'unsigned int' 1
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+ // CHECK-NEXT: gang clause static
+ // CHECK-NEXT: ImplicitCastExpr
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+ // CHECK-NEXT: gang clause static
+ // CHECK-NEXT: OpenACCAsteriskSizeExpr
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+ // CHECK-NEXT: gang clause dim
+ // CHECK-NEXT: ConstantExpr
+ // CHECK-NEXT: value: Int 1
+ // CHECK-NEXT: SubstNonTypeTemplateParmExpr
+ // CHECK-NEXT: NonTypeTemplateParmDecl
+ // CHECK-NEXT: IntegerLiteral{{.*}}'unsigned int' 1
+ // CHECK-NEXT: gang clause static
+ // CHECK-NEXT: ImplicitCastExpr
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+ // CHECK-NEXT: gang clause dim static
+ // CHECK-NEXT: ConstantExpr
+ // CHECK-NEXT: value: Int 1
+ // CHECK-NEXT: SubstNonTypeTemplateParmExpr
+ // CHECK-NEXT: NonTypeTemplateParmDecl
+ // CHECK-NEXT: IntegerLiteral{{.*}}'unsigned int' 1
+ // CHECK-NEXT: ImplicitCastExpr
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Val' 'int'
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+
+ // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} serial loop
+ // CHECK-NEXT: gang clause
+ // CHECK-NEXT: ForStmt
+ // CHECK: NullStmt
+}
+
+void inst() {
+ TemplateUses<int, 1>(5);
+}
+
+#endif // PCH_HELPER
diff --git a/clang/test/SemaOpenACC/combined-construct-gang-clause.cpp b/clang/test/SemaOpenACC/combined-construct-gang-clause.cpp
new file mode 100644
index 00000000000000..c5346c4cf5b90a
--- /dev/null
+++ b/clang/test/SemaOpenACC/combined-construct-gang-clause.cpp
@@ -0,0 +1,307 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+struct S{};
+struct Converts{
+ operator int();
+};
+
+template<typename T, unsigned Zero, unsigned Two, unsigned Four>
+void ParallelTempl() {
+ T i;
+
+ // expected-error at +1{{'num' argument on 'gang' clause is not permitted on a 'parallel loop' construct}}
+#pragma acc parallel loop gang(i)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{'num' argument on 'gang' clause is not permitted on a 'parallel loop' construct}}
+#pragma acc parallel loop gang(num:i)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{argument to 'gang' clause dimension must be a constant expression}}
+#pragma acc parallel loop gang(dim:i)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 0}}
+#pragma acc parallel loop gang(dim:Zero)
+ for(int i = 0; i < 5; ++i);
+
+#pragma acc parallel loop gang(dim:Two)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 4}}
+#pragma acc parallel loop gang(dim:Four)
+ for(int i = 0; i < 5; ++i);
+
+#pragma acc parallel loop gang(static:i) gang(dim:Two)
+ for(int i = 0; i < 5; ++i);
+
+#pragma acc parallel loop gang(static:i, dim:Two)
+ for(int i = 0; i < 5; ++i);
+
+#pragma acc parallel loop gang(dim:Two) gang(static:*)
+ for(int i = 0; i < 5; ++i);
+
+#pragma acc parallel loop gang(dim:Two, static:*)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +4{{OpenACC 'gang' clause may have at most one 'static' argument}}
+ // expected-note at +3{{previous expression is here}}
+ // expected-error at +2{{OpenACC 'gang' clause may have at most one 'dim' argument}}
+ // expected-note at +1{{previous expression is here}}
+#pragma acc parallel loop gang(dim:Two, static:*, dim:1, static:i)
+ for(int i = 0; i < 5; ++i);
+}
+
+void Parallel() {
+ ParallelTempl<int, 0, 2, 4>(); // expected-note{{in instantiation of function template}}
+ int i;
+
+ // expected-error at +1{{'num' argument on 'gang' clause is not permitted on a 'parallel loop' construct}}
+#pragma acc parallel loop gang(i)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{'num' argument on 'gang' clause is not permitted on a 'parallel loop' construct}}
+#pragma acc parallel loop gang(num:i)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{argument to 'gang' clause dimension must be a constant expression}}
+#pragma acc parallel loop gang(dim:i)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 0}}
+#pragma acc parallel loop gang(dim:0)
+ for(int i = 0; i < 5; ++i);
+
+#pragma acc parallel loop gang(dim:2)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 4}}
+#pragma acc parallel loop gang(dim:4)
+ for(int i = 0; i < 5; ++i);
+
+#pragma acc parallel loop gang(static:i) gang(dim:2)
+ for(int i = 0; i < 5; ++i);
+
+#pragma acc parallel loop gang(static:i, dim:2)
+ for(int i = 0; i < 5; ++i);
+
+#pragma acc parallel loop gang(dim:2) gang(static:*)
+ for(int i = 0; i < 5; ++i);
+
+#pragma acc parallel loop gang(dim:2, static:*)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +4{{OpenACC 'gang' clause may have at most one 'static' argument}}
+ // expected-note at +3{{previous expression is here}}
+ // expected-error at +2{{OpenACC 'gang' clause may have at most one 'dim' argument}}
+ // expected-note at +1{{previous expression is here}}
+#pragma acc parallel loop gang(dim:2, static:*, dim:1, static:i)
+ for(int i = 0; i < 5; ++i);
+}
+
+template<typename SomeS, typename SomeC, typename Int>
+void StaticIsIntegralTempl() {
+ SomeS s;
+ // expected-error at +1{{OpenACC clause 'gang' requires expression of integer type ('S' invalid)}}
+#pragma acc parallel loop gang(dim:2) gang(static:s)
+ for(int i = 0; i < 5; ++i);
+
+ SomeC C;
+#pragma acc parallel loop gang(dim:2) gang(static:C)
+ for(int i = 0; i < 5; ++i);
+ Int i;
+#pragma acc parallel loop gang(dim:2) gang(static:i)
+ for(int i = 0; i < 5; ++i);
+
+#pragma acc parallel loop gang(dim:2) gang(static:*)
+ for(int i = 0; i < 5; ++i);
+}
+
+void StaticIsIntegral() {
+ StaticIsIntegralTempl<S, Converts, int>();// expected-note{{in instantiation of function template}}
+
+ S s;
+ // expected-error at +1{{OpenACC clause 'gang' requires expression of integer type ('S' invalid)}}
+#pragma acc parallel loop gang(dim:2) gang(static:s)
+ for(int i = 0; i < 5; ++i);
+
+ Converts C;
+#pragma acc parallel loop gang(dim:2) gang(static:C)
+ for(int i = 0; i < 5; ++i);
+}
+
+template<unsigned I>
+void SerialTempl() {
+ // expected-error at +1{{'num' argument on 'gang' clause is not permitted on a 'serial loop'}}
+#pragma acc serial loop gang(I)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{'num' argument on 'gang' clause is not permitted on a 'serial loop'}}
+#pragma acc serial loop gang(num:I)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{'dim' argument on 'gang' clause is not permitted on a 'serial loop'}}
+#pragma acc serial loop gang(dim:I)
+ for(int i = 0; i < 5; ++i);
+
+#pragma acc serial loop gang(static:I)
+ for(int i = 0; i < 5; ++i);
+}
+
+void Serial() {
+ SerialTempl<2>();
+
+ // expected-error at +1{{'num' argument on 'gang' clause is not permitted on a 'serial loop'}}
+#pragma acc serial loop gang(1)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{'num' argument on 'gang' clause is not permitted on a 'serial loop'}}
+#pragma acc serial loop gang(num:1)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +1{{'dim' argument on 'gang' clause is not permitted on a 'serial loop'}}
+#pragma acc serial loop gang(dim:1)
+ for(int i = 0; i < 5; ++i);
+
+#pragma acc serial loop gang(static:1)
+ for(int i = 0; i < 5; ++i);
+
+ int i;
+#pragma acc serial loop gang(static:i)
+ for(int i = 0; i < 5; ++i);
+}
+
+template<typename T>
+void KernelsTempl() {
+ T t;
+ // expected-error at +1{{'dim' argument on 'gang' clause is not permitted on a 'kernels loop'}}
+#pragma acc kernels loop gang(dim:t)
+ for(int i = 0; i < 5; ++i);
+
+#pragma acc kernels loop gang(static:t)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +2{{'num_gangs' clause not allowed on a 'kernels loop' construct that has a 'gang' clause with a 'num' argument}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels loop gang(t) num_gangs(t)
+ for(int i = 0; i < 5; ++i);
+
+ // OK, kernels loop should block this.
+#pragma acc kernels num_gangs(t)
+#pragma acc kernels loop gang(static:t)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +2{{'num' argument to 'gang' clause not allowed on a 'kernels loop' construct that has a 'num_gangs' clause}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels loop num_gangs(t) gang(t)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +2{{'num_gangs' clause not allowed on a 'kernels loop' construct that has a 'gang' clause with a 'num' argument}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels loop gang(num:t) num_gangs(t)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +2{{'num' argument to 'gang' clause not allowed on a 'kernels loop' construct that has a 'num_gangs' clause}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels loop num_gangs(t) gang(num:t)
+ for(int i = 0; i < 5; ++i);
+}
+
+void Kernels() {
+ KernelsTempl<unsigned>();
+
+ // expected-error at +1{{'dim' argument on 'gang' clause is not permitted on a 'kernels loop'}}
+#pragma acc kernels loop gang(dim:1)
+ for(int i = 0; i < 5; ++i);
+
+ unsigned t;
+#pragma acc kernels loop gang(static:t)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +2{{'num_gangs' clause not allowed on a 'kernels loop' construct that has a 'gang' clause with a 'num' argument}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels loop gang(t) num_gangs(t)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +2{{'num' argument to 'gang' clause not allowed on a 'kernels loop' construct that has a 'num_gangs' clause}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels loop num_gangs(t) gang(t)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +2{{'num_gangs' clause not allowed on a 'kernels loop' construct that has a 'gang' clause with a 'num' argument}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels loop gang(num:t) num_gangs(t)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +2{{'num' argument to 'gang' clause not allowed on a 'kernels loop' construct that has a 'num_gangs' clause}}
+ // expected-note at +1{{previous clause is here}}
+#pragma acc kernels loop num_gangs(t) gang(num:t)
+ for(int i = 0; i < 5; ++i);
+
+ // OK, intervening compute/combined construct.
+#pragma acc kernels loop gang(num:1)
+ for(int i = 0; i < 5; ++i) {
+#pragma acc serial loop gang(static:1)
+ for(int i = 0; i < 5; ++i);
+ }
+#pragma acc kernels loop gang(num:1)
+ for(int i = 0; i < 5; ++i) {
+#pragma acc serial
+#pragma acc loop gang(static:1)
+ for(int i = 0; i < 5; ++i);
+ }
+#pragma acc kernels
+#pragma acc loop gang(num:1)
+ for(int i = 0; i < 5; ++i) {
+#pragma acc serial loop gang(static:1)
+ for(int i = 0; i < 5; ++i);
+ }
+
+#pragma acc kernels loop gang(num:1)
+ for(int i = 0; i < 5; ++i) {
+ // expected-error at +2{{loop with a 'gang' clause may not exist in the region of a 'gang' clause on a 'kernels loop' construct}}
+ // expected-note at -3{{previous clause is here}}
+#pragma acc loop gang(static:1)
+ for(int i = 0; i < 5; ++i);
+ }
+
+ // OK, on a
diff erent 'loop', not in the associated stmt.
+#pragma acc kernels loop gang(num:1)
+ for(int i = 0; i < 5; ++i);
+#pragma acc loop gang(static:1)
+ for(int i = 0; i < 5; ++i);
+
+ // OK, on a
diff erent 'loop', not in the associated stmt.
+#pragma acc kernels loop gang(num:1)
+ for(int i = 0; i < 5; ++i);
+#pragma acc kernels loop gang(static:1)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +2{{OpenACC 'gang' clause may have at most one unnamed or 'num' argument}}
+ // expected-note at +1{{previous expression is here}}
+#pragma acc kernels loop gang(5, num:1)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +2{{OpenACC 'gang' clause may have at most one unnamed or 'num' argument}}
+ // expected-note at +1{{previous expression is here}}
+#pragma acc kernels loop gang(num:5, 1)
+ for(int i = 0; i < 5; ++i);
+
+ // expected-error at +3{{OpenACC 'gang' clause may have at most one unnamed or 'num' argument}}
+ // expected-note at +2{{previous expression is here}}
+#pragma acc kernels
+#pragma acc kernels loop gang(num:5, num:1)
+ for(int i = 0; i < 5; ++i);
+}
+
+void MaxOneEntry() {
+ // expected-error at +2{{OpenACC 'gang' clause may have at most one 'static' argument}}
+ // expected-note at +1{{previous expression is here}}
+#pragma acc kernels loop gang(static: 1, static:1)
+ for(int i = 0; i < 5; ++i);
+
+#pragma acc kernels loop gang gang(static:1)
+ for(int i = 0; i < 5; ++i);
+}
+
+
diff --git a/clang/test/SemaOpenACC/loop-construct-gang-clause.cpp b/clang/test/SemaOpenACC/loop-construct-gang-clause.cpp
index a2bda5a7e82f63..6b4c03bbc50c6b 100644
--- a/clang/test/SemaOpenACC/loop-construct-gang-clause.cpp
+++ b/clang/test/SemaOpenACC/loop-construct-gang-clause.cpp
@@ -278,7 +278,7 @@ void Kernels() {
#pragma acc kernels
#pragma acc loop gang(num:1)
for(int j = 0; j < 5; ++j) {
- // expected-error at +2{{loop with a 'gang' clause may not exist in the region of a 'gang' clause on a 'kernels' compute construct}}
+ // expected-error at +2{{loop with a 'gang' clause may not exist in the region of a 'gang' clause on a 'kernels' construct}}
// expected-note at -3{{previous clause is here}}
#pragma acc loop gang(static:1)
for(int i = 0; i < 5; ++i);
More information about the cfe-commits
mailing list