[clang] 1b75b9e - [OpenACC] Handle sema for gang, worker, vector, seq clauses on routine
via cfe-commits
cfe-commits at lists.llvm.org
Thu Mar 6 11:53:51 PST 2025
Author: erichkeane
Date: 2025-03-06T11:53:46-08:00
New Revision: 1b75b9e665ee3c43de85c25f8d5f10d4efb3ca39
URL: https://github.com/llvm/llvm-project/commit/1b75b9e665ee3c43de85c25f8d5f10d4efb3ca39
DIFF: https://github.com/llvm/llvm-project/commit/1b75b9e665ee3c43de85c25f8d5f10d4efb3ca39.diff
LOG: [OpenACC] Handle sema for gang, worker, vector, seq clauses on routine
These 4 clauses are mutually exclusive, AND require at least one of
them. Additionally, gang has some additional restrictions in that only
the 'dim' specifier is permitted. This patch implements all of this, and
ends up refactoring the handling of each of these clauses for
readabililty.
Added:
clang/test/SemaOpenACC/routine-construct-clauses.cpp
Modified:
clang/include/clang/Sema/SemaOpenACC.h
clang/lib/AST/DeclPrinter.cpp
clang/lib/Parse/ParseOpenACC.cpp
clang/lib/Sema/SemaOpenACC.cpp
clang/lib/Sema/SemaOpenACCClause.cpp
clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
clang/lib/Serialization/ASTReader.cpp
clang/test/AST/ast-print-openacc-routine-construct.cpp
clang/test/ParserOpenACC/parse-clauses.c
clang/test/ParserOpenACC/parse-clauses.cpp
clang/test/ParserOpenACC/parse-constructs.c
clang/test/ParserOpenACC/parse-constructs.cpp
clang/test/SemaOpenACC/routine-construct-ast.cpp
clang/test/SemaOpenACC/routine-construct.cpp
clang/test/SemaOpenACC/unimplemented-construct.c
Removed:
################################################################################
diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h
index 442e53f903222..748dcdd251a92 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -698,7 +698,8 @@ class SemaOpenACC : public SemaBase {
/// parsing has consumed the 'annot_pragma_openacc_end' token. This DOES
/// happen before any associated declarations or statements have been parsed.
/// This function is only called when we are parsing a 'Decl' context.
- bool ActOnStartDeclDirective(OpenACCDirectiveKind K, SourceLocation StartLoc);
+ bool ActOnStartDeclDirective(OpenACCDirectiveKind K, SourceLocation StartLoc,
+ ArrayRef<const OpenACCClause *> Clauses);
/// Called when we encounter an associated statement for our construct, this
/// should check legality of the statement as it appertains to this Construct.
StmtResult ActOnAssociatedStmt(SourceLocation DirectiveLoc,
diff --git a/clang/lib/AST/DeclPrinter.cpp b/clang/lib/AST/DeclPrinter.cpp
index 959f7dc0d4669..0b59fa5aec1be 100644
--- a/clang/lib/AST/DeclPrinter.cpp
+++ b/clang/lib/AST/DeclPrinter.cpp
@@ -1918,9 +1918,12 @@ void DeclPrinter::VisitNonTypeTemplateParmDecl(
void DeclPrinter::VisitOpenACCDeclareDecl(OpenACCDeclareDecl *D) {
if (!D->isInvalidDecl()) {
- Out << "#pragma acc declare ";
- OpenACCClausePrinter Printer(Out, Policy);
- Printer.VisitClauseList(D->clauses());
+ Out << "#pragma acc declare";
+ if (!D->clauses().empty()) {
+ Out << ' ';
+ OpenACCClausePrinter Printer(Out, Policy);
+ Printer.VisitClauseList(D->clauses());
+ }
}
}
void DeclPrinter::VisitOpenACCRoutineDecl(OpenACCRoutineDecl *D) {
@@ -1941,7 +1944,10 @@ void DeclPrinter::VisitOpenACCRoutineDecl(OpenACCRoutineDecl *D) {
Out << ")";
}
- OpenACCClausePrinter Printer(Out, Policy);
- Printer.VisitClauseList(D->clauses());
+ if (!D->clauses().empty()) {
+ Out << ' ';
+ OpenACCClausePrinter Printer(Out, Policy);
+ Printer.VisitClauseList(D->clauses());
+ }
}
}
diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp
index 754bd26c448dd..6ea17c97d6345 100644
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@@ -1539,8 +1539,8 @@ Parser::DeclGroupPtrTy Parser::ParseOpenACCDirectiveDecl() {
OpenACCDirectiveParseInfo DirInfo = ParseOpenACCDirective();
- if (getActions().OpenACC().ActOnStartDeclDirective(DirInfo.DirKind,
- DirInfo.StartLoc))
+ if (getActions().OpenACC().ActOnStartDeclDirective(
+ DirInfo.DirKind, DirInfo.StartLoc, DirInfo.Clauses))
return nullptr;
return DeclGroupPtrTy::make(getActions().OpenACC().ActOnEndDeclDirective(
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index f4f105bf94e7e..4fe6bf5099a64 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -1555,7 +1555,7 @@ bool SemaOpenACC::ActOnStartStmtDirective(
// Declaration directives an appear in a statement location, so call into that
// function here.
if (K == OpenACCDirectiveKind::Declare || K == OpenACCDirectiveKind::Routine)
- return ActOnStartDeclDirective(K, StartLoc);
+ return ActOnStartDeclDirective(K, StartLoc, Clauses);
SemaRef.DiscardCleanupsInEvaluationContext();
SemaRef.PopExpressionEvaluationContext();
@@ -1832,14 +1832,29 @@ StmtResult SemaOpenACC::ActOnAssociatedStmt(
llvm_unreachable("Invalid associated statement application");
}
-bool SemaOpenACC::ActOnStartDeclDirective(OpenACCDirectiveKind K,
- SourceLocation StartLoc) {
+bool SemaOpenACC::ActOnStartDeclDirective(
+ OpenACCDirectiveKind K, SourceLocation StartLoc,
+ ArrayRef<const OpenACCClause *> Clauses) {
// OpenCC3.3 2.1 (line 889)
// A program must not depend on the order of evaluation of expressions in
// clause arguments or on any side effects of the evaluations.
SemaRef.DiscardCleanupsInEvaluationContext();
SemaRef.PopExpressionEvaluationContext();
+ if (K == OpenACCDirectiveKind::Routine &&
+ llvm::find_if(Clauses,
+ llvm::IsaPred<OpenACCGangClause, OpenACCWorkerClause,
+ OpenACCVectorClause, OpenACCSeqClause>) ==
+ Clauses.end())
+ return Diag(StartLoc, diag::err_acc_construct_one_clause_of)
+ << K
+ << GetListOfClauses({
+ OpenACCClauseKind::Gang,
+ OpenACCClauseKind::Worker,
+ OpenACCClauseKind::Vector,
+ OpenACCClauseKind::Seq,
+ });
+
return diagnoseConstructAppertainment(*this, K, StartLoc, /*IsStmt=*/false);
}
diff --git a/clang/lib/Sema/SemaOpenACCClause.cpp b/clang/lib/Sema/SemaOpenACCClause.cpp
index 532bc3e712b3a..149481d337b17 100644
--- a/clang/lib/Sema/SemaOpenACCClause.cpp
+++ b/clang/lib/Sema/SemaOpenACCClause.cpp
@@ -633,9 +633,18 @@ class SemaOpenACCClauseVisitor {
// OpenACC 3.3 2.9:
// A 'gang', 'worker', or 'vector' clause may not appear if a 'seq' clause
// appears.
- bool DiagIfSeqClause(SemaOpenACC::OpenACCParsedClause &Clause) {
+ // -also-
+ // OpenACC3.3 2.15: (routine)
+ // Exactly one of the 'gang', 'worker', 'vector' or 'seq' clauses must appear.
+ bool
+ DiagGangWorkerVectorSeqConflict(SemaOpenACC::OpenACCParsedClause &Clause) {
const auto *Itr =
- llvm::find_if(ExistingClauses, llvm::IsaPred<OpenACCSeqClause>);
+ Clause.getDirectiveKind() == OpenACCDirectiveKind::Routine
+ ? llvm::find_if(
+ ExistingClauses,
+ llvm::IsaPred<OpenACCSeqClause, OpenACCGangClause,
+ OpenACCWorkerClause, OpenACCVectorClause>)
+ : llvm::find_if(ExistingClauses, llvm::IsaPred<OpenACCSeqClause>);
if (Itr != ExistingClauses.end()) {
SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_cannot_combine)
@@ -1329,6 +1338,38 @@ ExprResult DiagIntArgInvalid(SemaOpenACC &S, Expr *E, StringRef TagKind,
return ExprError();
}
+ExprResult CheckGangDimExpr(SemaOpenACC &S, Expr *E) {
+ // 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.
+ // -also-
+ // OpenACC 3.3 2.15: The 'dim' argument must be a constant positive integer
+ // with 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})};
+}
+
ExprResult CheckGangParallelExpr(SemaOpenACC &S, OpenACCDirectiveKind DK,
OpenACCDirectiveKind AssocKind,
OpenACCGangKind GK, Expr *E) {
@@ -1340,35 +1381,8 @@ ExprResult CheckGangParallelExpr(SemaOpenACC &S, OpenACCDirectiveKind DK,
// 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})};
- }
+ case OpenACCGangKind::Dim:
+ return CheckGangDimExpr(S, E);
}
llvm_unreachable("Unknown gang kind in gang parallel check");
}
@@ -1433,21 +1447,32 @@ ExprResult CheckGangSerialExpr(SemaOpenACC &S, OpenACCDirectiveKind DK,
llvm_unreachable("Unknown gang kind in gang serial check");
}
+ExprResult CheckGangRoutineExpr(SemaOpenACC &S, OpenACCDirectiveKind DK,
+ OpenACCDirectiveKind AssocKind,
+ OpenACCGangKind GK, Expr *E) {
+ switch (GK) {
+ // Only 'dim' is allowed on a routine, so diallow num and static.
+ case OpenACCGangKind::Num:
+ case OpenACCGangKind::Static:
+ return DiagIntArgInvalid(S, E, GK, OpenACCClauseKind::Gang, DK, AssocKind);
+ case OpenACCGangKind::Dim:
+ return CheckGangDimExpr(S, E);
+ }
+ llvm_unreachable("Unknown gang kind in gang serial check");
+}
+
OpenACCClause *SemaOpenACCClauseVisitor::VisitVectorClause(
SemaOpenACC::OpenACCParsedClause &Clause) {
- if (DiagIfSeqClause(Clause))
+ if (DiagGangWorkerVectorSeqConflict(Clause))
return nullptr;
- // Restrictions only properly implemented on 'loop'/'combined' constructs, and
- // it is the only construct that can do anything with this, so skip/treat as
- // unimplemented for the routine constructs.
- if (!isDirectiveKindImplemented(Clause.getDirectiveKind()))
- return isNotImplemented();
-
Expr *IntExpr =
Clause.getNumIntExprs() != 0 ? Clause.getIntExprs()[0] : nullptr;
if (IntExpr) {
- if (!isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind())) {
+ switch (Clause.getDirectiveKind()) {
+ default:
+ llvm_unreachable("Invalid directive kind for this clause");
+ case OpenACCDirectiveKind::Loop:
switch (SemaRef.getActiveComputeConstructInfo().Kind) {
case OpenACCDirectiveKind::Invalid:
case OpenACCDirectiveKind::Parallel:
@@ -1483,34 +1508,38 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitVectorClause(
default:
llvm_unreachable("Non compute construct in active compute construct");
}
- } else {
- if (Clause.getDirectiveKind() == OpenACCDirectiveKind::SerialLoop) {
- DiagIntArgInvalid(SemaRef, IntExpr, "length", OpenACCClauseKind::Vector,
- Clause.getDirectiveKind(),
- SemaRef.getActiveComputeConstructInfo().Kind);
- IntExpr = nullptr;
- } else if (Clause.getDirectiveKind() ==
- OpenACCDirectiveKind::KernelsLoop) {
- const auto *Itr = llvm::find_if(
- ExistingClauses, llvm::IsaPred<OpenACCVectorLengthClause>);
- if (Itr != ExistingClauses.end()) {
- SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict)
- << "length" << 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);
+ break;
+ case OpenACCDirectiveKind::KernelsLoop: {
+ const auto *Itr = llvm::find_if(ExistingClauses,
+ llvm::IsaPred<OpenACCVectorLengthClause>);
+ if (Itr != ExistingClauses.end()) {
+ SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict)
+ << "length" << 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);
- IntExpr = nullptr;
- }
+ IntExpr = nullptr;
}
+ break;
+ }
+ case OpenACCDirectiveKind::SerialLoop:
+ case OpenACCDirectiveKind::Routine:
+ DiagIntArgInvalid(SemaRef, IntExpr, "length", OpenACCClauseKind::Vector,
+ Clause.getDirectiveKind(),
+ SemaRef.getActiveComputeConstructInfo().Kind);
+ IntExpr = nullptr;
+ break;
+ case OpenACCDirectiveKind::ParallelLoop:
+ break;
}
}
- if (!isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind())) {
+ if (Clause.getDirectiveKind() == OpenACCDirectiveKind::Loop) {
// 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.
@@ -1533,20 +1562,17 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitVectorClause(
OpenACCClause *SemaOpenACCClauseVisitor::VisitWorkerClause(
SemaOpenACC::OpenACCParsedClause &Clause) {
- if (DiagIfSeqClause(Clause))
+ if (DiagGangWorkerVectorSeqConflict(Clause))
return nullptr;
- // Restrictions only properly implemented on 'loop'/'combined' constructs, and
- // it is the only construct that can do anything with this, so skip/treat as
- // unimplemented for the routine constructs.
- if (!isDirectiveKindImplemented(Clause.getDirectiveKind()))
- return isNotImplemented();
-
Expr *IntExpr =
Clause.getNumIntExprs() != 0 ? Clause.getIntExprs()[0] : nullptr;
if (IntExpr) {
- if (!isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind())) {
+ switch (Clause.getDirectiveKind()) {
+ default:
+ llvm_unreachable("Invalid directive kind for this clause");
+ case OpenACCDirectiveKind::Loop:
switch (SemaRef.getActiveComputeConstructInfo().Kind) {
case OpenACCDirectiveKind::Invalid:
case OpenACCDirectiveKind::ParallelLoop:
@@ -1580,35 +1606,35 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitWorkerClause(
default:
llvm_unreachable("Non compute construct in active compute construct");
}
- } else {
- if (Clause.getDirectiveKind() == OpenACCDirectiveKind::ParallelLoop ||
- Clause.getDirectiveKind() == OpenACCDirectiveKind::SerialLoop) {
- DiagIntArgInvalid(SemaRef, IntExpr, OpenACCGangKind::Num,
- OpenACCClauseKind::Worker, Clause.getDirectiveKind(),
- SemaRef.getActiveComputeConstructInfo().Kind);
- IntExpr = nullptr;
- } else {
- assert(Clause.getDirectiveKind() == OpenACCDirectiveKind::KernelsLoop &&
- "Unknown combined directive kind?");
- const auto *Itr = llvm::find_if(ExistingClauses,
- llvm::IsaPred<OpenACCNumWorkersClause>);
- if (Itr != ExistingClauses.end()) {
- SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict)
- << "num" << 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);
+ break;
+ case OpenACCDirectiveKind::ParallelLoop:
+ case OpenACCDirectiveKind::SerialLoop:
+ case OpenACCDirectiveKind::Routine:
+ DiagIntArgInvalid(SemaRef, IntExpr, OpenACCGangKind::Num,
+ OpenACCClauseKind::Worker, Clause.getDirectiveKind(),
+ SemaRef.getActiveComputeConstructInfo().Kind);
+ IntExpr = nullptr;
+ break;
+ case OpenACCDirectiveKind::KernelsLoop: {
+ const auto *Itr = llvm::find_if(ExistingClauses,
+ llvm::IsaPred<OpenACCNumWorkersClause>);
+ if (Itr != ExistingClauses.end()) {
+ SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict)
+ << "num" << 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);
- IntExpr = nullptr;
- }
+ IntExpr = nullptr;
}
}
+ }
}
- if (!isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind())) {
+ if (Clause.getDirectiveKind() == OpenACCDirectiveKind::Loop) {
// 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.
@@ -1645,15 +1671,9 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitWorkerClause(
OpenACCClause *SemaOpenACCClauseVisitor::VisitGangClause(
SemaOpenACC::OpenACCParsedClause &Clause) {
- if (DiagIfSeqClause(Clause))
+ if (DiagGangWorkerVectorSeqConflict(Clause))
return nullptr;
- // 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 (!isDirectiveKindImplemented(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.
@@ -1721,7 +1741,7 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitGangClause(
IntExprs.push_back(ER.get());
}
- if (!isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind())) {
+ if (Clause.getDirectiveKind() == OpenACCDirectiveKind::Loop) {
// 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
@@ -1790,31 +1810,36 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitIfPresentClause(
OpenACCClause *SemaOpenACCClauseVisitor::VisitSeqClause(
SemaOpenACC::OpenACCParsedClause &Clause) {
- // Restrictions only properly implemented on 'loop' constructs and combined ,
- // and it is the only construct that can do anything with this, so skip/treat
- // as unimplemented for the routine constructs.
- if (!isDirectiveKindImplemented(Clause.getDirectiveKind()))
- return isNotImplemented();
- // OpenACC 3.3 2.9:
- // Only one of the seq, independent, and auto clauses may appear.
- const auto *Itr =
- llvm::find_if(ExistingClauses,
- llvm::IsaPred<OpenACCAutoClause, OpenACCIndependentClause,
- OpenACCSeqClause>);
- if (Itr != ExistingClauses.end()) {
- SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_loop_spec_conflict)
- << Clause.getClauseKind() << Clause.getDirectiveKind();
- SemaRef.Diag((*Itr)->getBeginLoc(), diag::note_acc_previous_clause_here);
- return nullptr;
+ if (Clause.getDirectiveKind() != OpenACCDirectiveKind::Routine) {
+ // OpenACC 3.3 2.9:
+ // Only one of the seq, independent, and auto clauses may appear.
+ const auto *Itr =
+ llvm::find_if(ExistingClauses,
+ llvm::IsaPred<OpenACCAutoClause, OpenACCIndependentClause,
+ OpenACCSeqClause>);
+ if (Itr != ExistingClauses.end()) {
+ SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_loop_spec_conflict)
+ << Clause.getClauseKind() << Clause.getDirectiveKind();
+ SemaRef.Diag((*Itr)->getBeginLoc(), diag::note_acc_previous_clause_here);
+ return nullptr;
+ }
}
// OpenACC 3.3 2.9:
// A 'gang', 'worker', or 'vector' clause may not appear if a 'seq' clause
// appears.
- Itr = llvm::find_if(ExistingClauses,
- llvm::IsaPred<OpenACCGangClause, OpenACCWorkerClause,
- OpenACCVectorClause>);
+ // -also-
+ // OpenACC3.3 2.15: (routine)
+ // Exactly one of the 'gang', 'worker', 'vector' or 'seq' clauses must appear.
+ const auto *Itr =
+ Clause.getDirectiveKind() == OpenACCDirectiveKind::Routine
+ ? llvm::find_if(ExistingClauses,
+ llvm::IsaPred<OpenACCGangClause, OpenACCWorkerClause,
+ OpenACCVectorClause, OpenACCSeqClause>)
+ : llvm::find_if(ExistingClauses,
+ llvm::IsaPred<OpenACCGangClause, OpenACCWorkerClause,
+ OpenACCVectorClause>);
if (Itr != ExistingClauses.end()) {
SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_cannot_combine)
@@ -2199,6 +2224,9 @@ SemaOpenACC::CheckGangExpr(ArrayRef<const OpenACCClause *> ExistingClauses,
case OpenACCDirectiveKind::KernelsLoop:
return CheckGangKernelsExpr(*this, ExistingClauses, DK,
ActiveComputeConstructInfo.Kind, GK, E);
+ case OpenACCDirectiveKind::Routine:
+ return CheckGangRoutineExpr(*this, DK, ActiveComputeConstructInfo.Kind, GK,
+ E);
case OpenACCDirectiveKind::Loop:
switch (ActiveComputeConstructInfo.Kind) {
case OpenACCDirectiveKind::Invalid:
@@ -2218,8 +2246,6 @@ SemaOpenACC::CheckGangExpr(ArrayRef<const OpenACCClause *> ExistingClauses,
llvm_unreachable("Non compute construct in active compute construct?");
}
default:
- // 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?");
@@ -2231,31 +2257,34 @@ SemaOpenACC::CheckGangClause(OpenACCDirectiveKind DirKind,
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 << DirKind;
- Diag((*ReductionItr)->getBeginLoc(),
- diag::note_acc_previous_clause_here);
- return nullptr;
+ // Reduction isn't possible on 'routine' so we don't bother checking it here.
+ if (DirKind != OpenACCDirectiveKind::Routine) {
+ // 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 << DirKind;
+ Diag((*ReductionItr)->getBeginLoc(),
+ diag::note_acc_previous_clause_here);
+ return nullptr;
+ }
}
}
}
diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
index 7294e42214bea..79075c23c3ef7 100644
--- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
+++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
@@ -1004,14 +1004,17 @@ class OpenACCDeclClauseInstantiator final
: public OpenACCClauseVisitor<OpenACCDeclClauseInstantiator> {
Sema &SemaRef;
const MultiLevelTemplateArgumentList &MLTAL;
+ ArrayRef<OpenACCClause *> ExistingClauses;
SemaOpenACC::OpenACCParsedClause &ParsedClause;
OpenACCClause *NewClause = nullptr;
public:
OpenACCDeclClauseInstantiator(Sema &S,
const MultiLevelTemplateArgumentList &MLTAL,
+ ArrayRef<OpenACCClause *> ExistingClauses,
SemaOpenACC::OpenACCParsedClause &ParsedClause)
- : SemaRef(S), MLTAL(MLTAL), ParsedClause(ParsedClause) {}
+ : SemaRef(S), MLTAL(MLTAL), ExistingClauses(ExistingClauses),
+ ParsedClause(ParsedClause) {}
OpenACCClause *CreatedClause() { return NewClause; }
#define VISIT_CLAUSE(CLAUSE_NAME) \
@@ -1056,7 +1059,6 @@ CLAUSE_NOT_ON_DECLS(DeviceNum)
CLAUSE_NOT_ON_DECLS(DeviceType)
CLAUSE_NOT_ON_DECLS(Finalize)
CLAUSE_NOT_ON_DECLS(FirstPrivate)
-CLAUSE_NOT_ON_DECLS(Gang)
CLAUSE_NOT_ON_DECLS(Host)
CLAUSE_NOT_ON_DECLS(If)
CLAUSE_NOT_ON_DECLS(IfPresent)
@@ -1067,15 +1069,63 @@ CLAUSE_NOT_ON_DECLS(NumWorkers)
CLAUSE_NOT_ON_DECLS(Private)
CLAUSE_NOT_ON_DECLS(Reduction)
CLAUSE_NOT_ON_DECLS(Self)
-CLAUSE_NOT_ON_DECLS(Seq)
CLAUSE_NOT_ON_DECLS(Tile)
CLAUSE_NOT_ON_DECLS(UseDevice)
-CLAUSE_NOT_ON_DECLS(Vector)
CLAUSE_NOT_ON_DECLS(VectorLength)
CLAUSE_NOT_ON_DECLS(Wait)
-CLAUSE_NOT_ON_DECLS(Worker)
#undef CLAUSE_NOT_ON_DECLS
+void OpenACCDeclClauseInstantiator::VisitGangClause(
+ const OpenACCGangClause &C) {
+ llvm::SmallVector<OpenACCGangKind> TransformedGangKinds;
+ llvm::SmallVector<Expr *> TransformedIntExprs;
+ assert(C.getNumExprs() <= 1 &&
+ "Only 1 expression allowed on gang clause in routine");
+
+ if (C.getNumExprs() > 0) {
+ assert(C.getExpr(0).first == OpenACCGangKind::Dim &&
+ "Only dim allowed on routine");
+ ExprResult ER =
+ SemaRef.SubstExpr(const_cast<Expr *>(C.getExpr(0).second), MLTAL);
+ if (ER.isUsable()) {
+ ER = SemaRef.OpenACC().CheckGangExpr(ExistingClauses,
+ ParsedClause.getDirectiveKind(),
+ C.getExpr(0).first, ER.get());
+ if (ER.isUsable()) {
+ TransformedGangKinds.push_back(OpenACCGangKind::Dim);
+ TransformedIntExprs.push_back(ER.get());
+ }
+ }
+ }
+
+ NewClause = SemaRef.OpenACC().CheckGangClause(
+ ParsedClause.getDirectiveKind(), ExistingClauses,
+ ParsedClause.getBeginLoc(), ParsedClause.getLParenLoc(),
+ TransformedGangKinds, TransformedIntExprs, ParsedClause.getEndLoc());
+}
+
+void OpenACCDeclClauseInstantiator::VisitSeqClause(const OpenACCSeqClause &C) {
+ NewClause = OpenACCSeqClause::Create(SemaRef.getASTContext(),
+ ParsedClause.getBeginLoc(),
+ ParsedClause.getEndLoc());
+}
+
+void OpenACCDeclClauseInstantiator::VisitWorkerClause(
+ const OpenACCWorkerClause &C) {
+ assert(!C.hasIntExpr() && "Int Expr not allowed on routine 'worker' clause");
+ NewClause = OpenACCWorkerClause::Create(SemaRef.getASTContext(),
+ ParsedClause.getBeginLoc(), {},
+ nullptr, ParsedClause.getEndLoc());
+}
+
+void OpenACCDeclClauseInstantiator::VisitVectorClause(
+ const OpenACCVectorClause &C) {
+ assert(!C.hasIntExpr() && "Int Expr not allowed on routine 'vector' clause");
+ NewClause = OpenACCVectorClause::Create(SemaRef.getASTContext(),
+ ParsedClause.getBeginLoc(), {},
+ nullptr, ParsedClause.getEndLoc());
+}
+
void OpenACCDeclClauseInstantiator::VisitCopyClause(
const OpenACCCopyClause &C) {
ParsedClause.setVarListDetails(VisitVarList(C.getVarList()),
@@ -1197,7 +1247,8 @@ llvm::SmallVector<OpenACCClause *> InstantiateOpenACCClauseList(
if (const auto *WithParms = dyn_cast<OpenACCClauseWithParams>(Clause))
ParsedClause.setLParenLoc(WithParms->getLParenLoc());
- OpenACCDeclClauseInstantiator Instantiator{S, MLTAL, ParsedClause};
+ OpenACCDeclClauseInstantiator Instantiator{S, MLTAL, TransformedClauses,
+ ParsedClause};
Instantiator.Visit(Clause);
if (Instantiator.CreatedClause())
TransformedClauses.push_back(Instantiator.CreatedClause());
@@ -1213,8 +1264,8 @@ Decl *TemplateDeclInstantiator::VisitOpenACCDeclareDecl(OpenACCDeclareDecl *D) {
InstantiateOpenACCClauseList(SemaRef, TemplateArgs, D->getDirectiveKind(),
D->clauses());
- if (SemaRef.OpenACC().ActOnStartDeclDirective(D->getDirectiveKind(),
- D->getBeginLoc()))
+ if (SemaRef.OpenACC().ActOnStartDeclDirective(
+ D->getDirectiveKind(), D->getBeginLoc(), TransformedClauses))
return nullptr;
DeclGroupRef Res = SemaRef.OpenACC().ActOnEndDeclDirective(
@@ -1242,8 +1293,8 @@ Decl *TemplateDeclInstantiator::VisitOpenACCRoutineDecl(OpenACCRoutineDecl *D) {
// the function decl is empty.
}
- if (SemaRef.OpenACC().ActOnStartDeclDirective(D->getDirectiveKind(),
- D->getBeginLoc()))
+ if (SemaRef.OpenACC().ActOnStartDeclDirective(
+ D->getDirectiveKind(), D->getBeginLoc(), TransformedClauses))
return nullptr;
DeclGroupRef Res = SemaRef.OpenACC().ActOnEndDeclDirective(
diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
index ca09c3d79d941..2c186e4047dde 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -12764,7 +12764,8 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
llvm::SmallVector<Expr *> Exprs;
for (unsigned I = 0; I < NumExprs; ++I) {
GangKinds.push_back(readEnum<OpenACCGangKind>());
- Exprs.push_back(readSubExpr());
+ // Can't use `readSubExpr` because this is usable from a 'decl' construct.
+ Exprs.push_back(readExpr());
}
return OpenACCGangClause::Create(getContext(), BeginLoc, LParenLoc,
GangKinds, Exprs, EndLoc);
diff --git a/clang/test/AST/ast-print-openacc-routine-construct.cpp b/clang/test/AST/ast-print-openacc-routine-construct.cpp
index 64628d5c64cbc..c199db71a4e48 100644
--- a/clang/test/AST/ast-print-openacc-routine-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-routine-construct.cpp
@@ -1,89 +1,91 @@
// RUN: %clang_cc1 -fopenacc -ast-print %s -o - | FileCheck %s
auto Lambda = [](){};
-// CHECK: #pragma acc routine(Lambda)
-#pragma acc routine(Lambda)
+// CHECK: #pragma acc routine(Lambda) worker
+#pragma acc routine(Lambda) worker
int function();
-// CHECK: #pragma acc routine(function)
-#pragma acc routine (function)
+// CHECK: #pragma acc routine(function) vector
+#pragma acc routine (function) vector
namespace NS {
int NSFunc();
auto Lambda = [](){};
}
-// CHECK: #pragma acc routine(NS::NSFunc)
-#pragma acc routine(NS::NSFunc)
-// CHECK: #pragma acc routine(NS::Lambda)
-#pragma acc routine(NS::Lambda)
+// CHECK: #pragma acc routine(NS::NSFunc) seq
+#pragma acc routine(NS::NSFunc) seq
+// CHECK: #pragma acc routine(NS::Lambda) gang
+#pragma acc routine(NS::Lambda) gang
+
+constexpr int getInt() { return 1; }
struct S {
void MemFunc();
static void StaticMemFunc();
constexpr static auto Lambda = [](){};
-// CHECK: #pragma acc routine(S::MemFunc)
-#pragma acc routine(S::MemFunc)
-// CHECK: #pragma acc routine(S::StaticMemFunc)
-#pragma acc routine(S::StaticMemFunc)
-// CHECK: #pragma acc routine(S::Lambda)
-#pragma acc routine(S::Lambda)
+// CHECK: #pragma acc routine(S::MemFunc) gang(dim: 1)
+#pragma acc routine(S::MemFunc) gang(dim:1)
+// CHECK: #pragma acc routine(S::StaticMemFunc) gang(dim: getInt())
+#pragma acc routine(S::StaticMemFunc) gang(dim:getInt())
+// CHECK: #pragma acc routine(S::Lambda) worker
+#pragma acc routine(S::Lambda) worker
-// CHECK: #pragma acc routine(MemFunc)
-#pragma acc routine(MemFunc)
-// CHECK: #pragma acc routine(StaticMemFunc)
-#pragma acc routine(StaticMemFunc)
-// CHECK: #pragma acc routine(Lambda)
-#pragma acc routine(Lambda)
+// CHECK: #pragma acc routine(MemFunc) gang(dim: 1)
+#pragma acc routine(MemFunc) gang(dim:1)
+// CHECK: #pragma acc routine(StaticMemFunc) gang(dim: getInt())
+#pragma acc routine(StaticMemFunc) gang(dim:getInt())
+// CHECK: #pragma acc routine(Lambda) worker
+#pragma acc routine(Lambda) worker
};
-// CHECK: #pragma acc routine(S::MemFunc)
-#pragma acc routine(S::MemFunc)
-// CHECK: #pragma acc routine(S::StaticMemFunc)
-#pragma acc routine(S::StaticMemFunc)
-// CHECK: #pragma acc routine(S::Lambda)
-#pragma acc routine(S::Lambda)
+// CHECK: #pragma acc routine(S::MemFunc) gang(dim: 1)
+#pragma acc routine(S::MemFunc) gang(dim:1)
+// CHECK: #pragma acc routine(S::StaticMemFunc) worker
+#pragma acc routine(S::StaticMemFunc) worker
+// CHECK: #pragma acc routine(S::Lambda) vector
+#pragma acc routine(S::Lambda) vector
template<typename T>
struct DepS {
void MemFunc();
static void StaticMemFunc();
- constexpr static auto Lambda = [](){};
+ constexpr static auto Lambda = [](){ return 1;};
-// CHECK: #pragma acc routine(Lambda)
-#pragma acc routine(Lambda)
-// CHECK: #pragma acc routine(MemFunc)
-#pragma acc routine(MemFunc)
-// CHECK: #pragma acc routine(StaticMemFunc)
-#pragma acc routine(StaticMemFunc)
+// CHECK: #pragma acc routine(Lambda) gang(dim: Lambda())
+#pragma acc routine(Lambda) gang(dim:Lambda())
+// CHECK: #pragma acc routine(MemFunc) worker
+#pragma acc routine(MemFunc) worker
+// CHECK: #pragma acc routine(StaticMemFunc) seq
+#pragma acc routine(StaticMemFunc) seq
-// CHECK: #pragma acc routine(DepS<T>::Lambda)
-#pragma acc routine(DepS::Lambda)
-// CHECK: #pragma acc routine(DepS<T>::MemFunc)
-#pragma acc routine(DepS::MemFunc)
-// CHECK: #pragma acc routine(DepS<T>::StaticMemFunc)
-#pragma acc routine(DepS::StaticMemFunc)
+// CHECK: #pragma acc routine(DepS<T>::Lambda) gang(dim: 1)
+#pragma acc routine(DepS::Lambda) gang(dim:1)
+// CHECK: #pragma acc routine(DepS<T>::MemFunc) gang
+#pragma acc routine(DepS::MemFunc) gang
+// CHECK: #pragma acc routine(DepS<T>::StaticMemFunc) worker
+#pragma acc routine(DepS::StaticMemFunc) worker
-// CHECK: #pragma acc routine(DepS<T>::Lambda)
-#pragma acc routine(DepS<T>::Lambda)
-// CHECK: #pragma acc routine(DepS<T>::MemFunc)
-#pragma acc routine(DepS<T>::MemFunc)
-// CHECK: #pragma acc routine(DepS<T>::StaticMemFunc)
-#pragma acc routine(DepS<T>::StaticMemFunc)
+// CHECK: #pragma acc routine(DepS<T>::Lambda) vector
+#pragma acc routine(DepS<T>::Lambda) vector
+// CHECK: #pragma acc routine(DepS<T>::MemFunc) seq
+#pragma acc routine(DepS<T>::MemFunc) seq
+// CHECK: #pragma acc routine(DepS<T>::StaticMemFunc) worker
+#pragma acc routine(DepS<T>::StaticMemFunc) worker
};
-// CHECK: #pragma acc routine(DepS<int>::Lambda)
-#pragma acc routine(DepS<int>::Lambda)
-// CHECK: #pragma acc routine(DepS<int>::MemFunc)
-#pragma acc routine(DepS<int>::MemFunc)
-// CHECK: #pragma acc routine(DepS<int>::StaticMemFunc)
-#pragma acc routine(DepS<int>::StaticMemFunc)
+// CHECK: #pragma acc routine(DepS<int>::Lambda) gang
+#pragma acc routine(DepS<int>::Lambda) gang
+// CHECK: #pragma acc routine(DepS<int>::MemFunc) gang(dim: 1)
+#pragma acc routine(DepS<int>::MemFunc) gang(dim:1)
+// CHECK: #pragma acc routine(DepS<int>::StaticMemFunc) vector
+#pragma acc routine(DepS<int>::StaticMemFunc) vector
template<typename T>
void TemplFunc() {
-// CHECK: #pragma acc routine(T::MemFunc)
-#pragma acc routine(T::MemFunc)
-// CHECK: #pragma acc routine(T::StaticMemFunc)
-#pragma acc routine(T::StaticMemFunc)
-// CHECK: #pragma acc routine(T::Lambda)
-#pragma acc routine(T::Lambda)
+// CHECK: #pragma acc routine(T::MemFunc) gang(dim: T::SomethingElse())
+#pragma acc routine(T::MemFunc) gang(dim:T::SomethingElse())
+// CHECK: #pragma acc routine(T::StaticMemFunc) worker
+#pragma acc routine(T::StaticMemFunc) worker
+// CHECK: #pragma acc routine(T::Lambda) seq
+#pragma acc routine(T::Lambda) seq
}
diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c
index ad894c3bcd04d..6b1bb208704d7 100644
--- a/clang/test/ParserOpenACC/parse-clauses.c
+++ b/clang/test/ParserOpenACC/parse-clauses.c
@@ -1261,18 +1261,20 @@ void Gang() {
}
- // expected-warning at +5{{OpenACC clause 'worker' not yet implemented, clause ignored}}
- // expected-warning at +4{{OpenACC clause 'vector' not yet implemented, clause ignored}}
- // expected-warning at +3{{OpenACC clause 'seq' not yet implemented, clause ignored}}
- // expected-warning at +2{{OpenACC clause 'nohost' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}}
+ // expected-warning at +6{{OpenACC clause 'nohost' not yet implemented, clause ignored}}
+ // expected-warning at +5{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}}
+ // expected-error at +4{{OpenACC clause 'seq' may not appear on the same construct as a 'worker' clause on a 'routine' construct}}
+ // expected-note at +3{{previous clause is here}}
+ // expected-error at +2{{OpenACC clause 'vector' may not appear on the same construct as a 'worker' clause on a 'routine' construct}}
+ // expected-note at +1{{previous clause is here}}
#pragma acc routine worker, vector, seq, nohost
void bar();
- // expected-warning at +4{{OpenACC clause 'worker' not yet implemented, clause ignored}}
- // expected-warning at +3{{OpenACC clause 'vector' not yet implemented, clause ignored}}
- // expected-warning at +2{{OpenACC clause 'seq' not yet implemented, clause ignored}}
- // expected-warning at +1{{OpenACC clause 'nohost' not yet implemented, clause ignored}}
+ // expected-warning at +5{{OpenACC clause 'nohost' not yet implemented, clause ignored}}
+ // expected-error at +4{{OpenACC clause 'seq' may not appear on the same construct as a 'worker' clause on a 'routine' construct}}
+ // expected-note at +3{{previous clause is here}}
+ // expected-error at +2{{OpenACC clause 'vector' may not appear on the same construct as a 'worker' clause on a 'routine' construct}}
+ // expected-note at +1{{previous clause is here}}
#pragma acc routine(bar) worker, vector, seq, nohost
@@ -1280,19 +1282,19 @@ void bar();
// expected-error at +2{{expected '('}}
// expected-warning at +1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}}
-#pragma acc routine bind
+#pragma acc routine seq bind
void BCP1();
// expected-error at +1{{expected identifier or string literal}}
-#pragma acc routine(BCP1) bind()
+#pragma acc routine(BCP1) seq bind()
// expected-warning at +2{{OpenACC clause 'bind' not yet implemented, clause ignored}}
// expected-warning at +1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}}
-#pragma acc routine bind("ReductionClauseParsing")
+#pragma acc routine seq bind("ReductionClauseParsing")
void BCP2();
// expected-warning at +1{{OpenACC clause 'bind' not yet implemented, clause ignored}}
-#pragma acc routine(BCP1) bind(BCP2)
+#pragma acc routine(BCP1) seq bind(BCP2)
// expected-error at +1{{use of undeclared identifier 'unknown_thing'}}
-#pragma acc routine(BCP1) bind(unknown_thing)
+#pragma acc routine(BCP1) seq bind(unknown_thing)
diff --git a/clang/test/ParserOpenACC/parse-clauses.cpp b/clang/test/ParserOpenACC/parse-clauses.cpp
index ff747f345cbb9..03a0fd945e0b9 100644
--- a/clang/test/ParserOpenACC/parse-clauses.cpp
+++ b/clang/test/ParserOpenACC/parse-clauses.cpp
@@ -66,35 +66,35 @@ void function();
// expected-warning at +1{{OpenACC clause 'bind' not yet implemented, clause ignored}}
-#pragma acc routine(use) bind(NS::NSFunc)
+#pragma acc routine(use) seq bind(NS::NSFunc)
// expected-error at +2{{'RecordTy' does not refer to a value}}
// expected-note@#RecTy{{declared here}}
-#pragma acc routine(use) bind(NS::RecordTy)
+#pragma acc routine(use) seq bind(NS::RecordTy)
// expected-error at +3{{'Value' is a private member of 'NS::RecordTy'}}
// expected-note@#VAL{{implicitly declared private here}}
// expected-warning at +1{{OpenACC clause 'bind' not yet implemented, clause ignored}}
-#pragma acc routine(use) bind(NS::RecordTy::Value)
+#pragma acc routine(use) seq bind(NS::RecordTy::Value)
// expected-warning at +1{{OpenACC clause 'bind' not yet implemented, clause ignored}}
-#pragma acc routine(use) bind(NS::RecordTy::ValuePub)
+#pragma acc routine(use) seq bind(NS::RecordTy::ValuePub)
// expected-warning at +1{{OpenACC clause 'bind' not yet implemented, clause ignored}}
-#pragma acc routine(use) bind(NS::TemplTy<int>)
+#pragma acc routine(use) seq bind(NS::TemplTy<int>)
// expected-error at +1{{no member named 'unknown' in namespace 'NS'}}
-#pragma acc routine(use) bind(NS::unknown<int>)
+#pragma acc routine(use) seq bind(NS::unknown<int>)
// expected-warning at +1{{OpenACC clause 'bind' not yet implemented, clause ignored}}
-#pragma acc routine(use) bind(NS::function)
+#pragma acc routine(use) seq bind(NS::function)
// expected-error at +3{{'priv_mem_function' is a private member of 'NS::RecordTy'}}
// expected-note@#PrivMemFun{{implicitly declared private here}}
// expected-warning at +1{{OpenACC clause 'bind' not yet implemented, clause ignored}}
-#pragma acc routine(use) bind(NS::RecordTy::priv_mem_function)
+#pragma acc routine(use) seq bind(NS::RecordTy::priv_mem_function)
// expected-warning at +1{{OpenACC clause 'bind' not yet implemented, clause ignored}}
-#pragma acc routine(use) bind(NS::RecordTy::mem_function)
+#pragma acc routine(use) seq bind(NS::RecordTy::mem_function)
// expected-error at +1{{string literal with user-defined suffix cannot be used here}}
-#pragma acc routine(use) bind("unknown udl"_UDL)
+#pragma acc routine(use) seq bind("unknown udl"_UDL)
// expected-warning at +2{{encoding prefix 'u' on an unevaluated string literal has no effect}}
// expected-warning at +1{{OpenACC clause 'bind' not yet implemented, clause ignored}}
-#pragma acc routine(use) bind(u"16 bits")
+#pragma acc routine(use) seq bind(u"16 bits")
// expected-warning at +2{{encoding prefix 'U' on an unevaluated string literal has no effect}}
// expected-warning at +1{{OpenACC clause 'bind' not yet implemented, clause ignored}}
-#pragma acc routine(use) bind(U"32 bits")
+#pragma acc routine(use) seq bind(U"32 bits")
diff --git a/clang/test/ParserOpenACC/parse-constructs.c b/clang/test/ParserOpenACC/parse-constructs.c
index 107d48a0b22d0..969c31b5bd8c1 100644
--- a/clang/test/ParserOpenACC/parse-constructs.c
+++ b/clang/test/ParserOpenACC/parse-constructs.c
@@ -151,29 +151,31 @@ void func() {
}
// expected-warning at +1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}}
-#pragma acc routine
+#pragma acc routine seq
void routine_func();
// expected-error at +2{{invalid OpenACC clause 'clause'}}
-// expected-warning at +1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}}
+// expected-error at +1{{OpenACC 'routine' construct must have at least one 'gang', 'worker', 'vector' or 'seq' clause}}
#pragma acc routine clause list
void routine_func();
// expected-error at +1{{use of undeclared identifier 'func_name'}}
-#pragma acc routine (func_name)
-// expected-error at +2{{use of undeclared identifier 'func_name'}}
-// expected-error at +1{{invalid OpenACC clause 'clause'}}
+#pragma acc routine (func_name) seq
+// expected-error at +3{{use of undeclared identifier 'func_name'}}
+// expected-error at +2{{invalid OpenACC clause 'clause'}}
+// expected-error at +1{{OpenACC 'routine' construct must have at least one 'gang', 'worker', 'vector' or 'seq' clause}}
#pragma acc routine (func_name) clause list
-#pragma acc routine (routine_func)
-// expected-error at +1{{invalid OpenACC clause 'clause'}}
+#pragma acc routine (routine_func) seq
+// expected-error at +2{{invalid OpenACC clause 'clause'}}
+// expected-error at +1{{OpenACC 'routine' construct must have at least one 'gang', 'worker', 'vector' or 'seq' clause}}
#pragma acc routine (routine_func) clause list
// expected-error at +2{{expected ')'}}
// expected-note at +1{{to match this '('}}
-#pragma acc routine (routine_func())
+#pragma acc routine (routine_func()) seq
// expected-error at +1{{expected identifier}}
-#pragma acc routine()
+#pragma acc routine() seq
// expected-error at +1{{expected identifier}}
-#pragma acc routine(int)
+#pragma acc routine(int) seq
diff --git a/clang/test/ParserOpenACC/parse-constructs.cpp b/clang/test/ParserOpenACC/parse-constructs.cpp
index 7892fb22fa2c2..2c083f558ab12 100644
--- a/clang/test/ParserOpenACC/parse-constructs.cpp
+++ b/clang/test/ParserOpenACC/parse-constructs.cpp
@@ -14,40 +14,40 @@ namespace NS {
}
// expected-error at +1{{use of undeclared identifier 'foo'; did you mean 'NS::foo'?}}
-#pragma acc routine(foo)
-#pragma acc routine(NS::foo)
+#pragma acc routine(foo) seq
+#pragma acc routine(NS::foo) seq
// expected-error at +2{{use of undeclared identifier 'templ'; did you mean 'NS::templ'?}}
// expected-error at +1{{OpenACC routine name 'NS::templ' names a set of overloads}}
-#pragma acc routine(templ)
+#pragma acc routine(templ) seq
// expected-error at +1{{OpenACC routine name 'NS::templ' names a set of overloads}}
-#pragma acc routine(NS::templ)
+#pragma acc routine(NS::templ) seq
// expected-error at +2{{use of undeclared identifier 'templ'; did you mean 'NS::templ'?}}
// expected-error at +1{{OpenACC routine name 'NS::templ' names a set of overloads}}
-#pragma acc routine(templ<int>)
+#pragma acc routine(templ<int>) seq
// expected-error at +1{{OpenACC routine name 'NS::templ<int>' names a set of overloads}}
-#pragma acc routine(NS::templ<int>)
+#pragma acc routine(NS::templ<int>) seq
// expected-error at +1{{use of undeclared identifier 'T'}}
-#pragma acc routine(templ<T>)
+#pragma acc routine(templ<T>) seq
// expected-error at +1{{use of undeclared identifier 'T'}}
-#pragma acc routine(NS::templ<T>)
+#pragma acc routine(NS::templ<T>) seq
// expected-error at +2{{expected ')'}}
// expected-note at +1{{to match this '('}}
-#pragma acc routine (NS::foo())
+#pragma acc routine (NS::foo()) seq
// expected-error at +1 {{expected unqualified-id}}
-#pragma acc routine()
+#pragma acc routine() seq
// expected-error at +1 {{expected unqualified-id}}
-#pragma acc routine(int)
+#pragma acc routine(int) seq
// expected-error at +2{{'C' does not refer to a value}}
// expected-note@#CDef{{declared here}}
-#pragma acc routine (NS::C)
+#pragma acc routine (NS::C) seq
// expected-error at +2{{'private_mem_func' is a private member of 'NS::C'}}
// expected-note@#PrivateMemFunc{{implicitly declared private here}}
-#pragma acc routine (NS::C::private_mem_func)
-#pragma acc routine (NS::C::public_mem_func)
+#pragma acc routine (NS::C::private_mem_func) seq
+#pragma acc routine (NS::C::public_mem_func) seq
diff --git a/clang/test/SemaOpenACC/routine-construct-ast.cpp b/clang/test/SemaOpenACC/routine-construct-ast.cpp
index 1ef54c9170565..f59317cf9d93a 100644
--- a/clang/test/SemaOpenACC/routine-construct-ast.cpp
+++ b/clang/test/SemaOpenACC/routine-construct-ast.cpp
@@ -7,110 +7,147 @@
#ifndef PCH_HELPER
#define PCH_HELPER
auto Lambda = [](){};
-#pragma acc routine(Lambda)
+#pragma acc routine(Lambda) worker
// CHECK: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' '(lambda at
+// CHECK-NEXT: worker clause
int function();
-#pragma acc routine (function)
+#pragma acc routine (function) vector
// CHECK: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'function' 'int ()'
+// CHECK-NEXT: vector clause
namespace NS {
// CHECK-NEXT: NamespaceDecl
int NSFunc();
auto Lambda = [](){};
}
-#pragma acc routine(NS::NSFunc)
+#pragma acc routine(NS::NSFunc) seq
// CHECK: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'NSFunc' 'int ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'NS'
-#pragma acc routine(NS::Lambda)
+// CHECK-NEXT: seq clause
+#pragma acc routine(NS::Lambda) gang
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'NS::(lambda at
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'NS'
+// CHECK-NEXT: gang clause
+
+constexpr int getInt() { return 1; }
struct S {
void MemFunc();
static void StaticMemFunc();
- constexpr static auto Lambda = [](){};
-#pragma acc routine(S::MemFunc)
+ constexpr static auto Lambda = [](){ return 1; };
+#pragma acc routine(S::MemFunc) gang(dim: 1)
// CHECK: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'void ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'S'
-#pragma acc routine(S::StaticMemFunc)
+// CHECK-NEXT: gang clause
+// CHECK-NEXT: ConstantExpr{{.*}}'int'
+// CHECK-NEXT: value: Int 1
+#pragma acc routine(S::StaticMemFunc) gang(dim:getInt())
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'void ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'S'
-#pragma acc routine(S::Lambda)
+// CHECK-NEXT: gang clause
+// CHECK-NEXT: ConstantExpr{{.*}}'int'
+// CHECK-NEXT: value: Int 1
+#pragma acc routine(S::Lambda) worker
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const S::(lambda at
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'S'
+// CHECK-NEXT: worker clause
-#pragma acc routine(MemFunc)
-// CHECK: OpenACCRoutineDecl{{.*}} routine name_specified
+#pragma acc routine(MemFunc) gang(dim: 1)
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'void ()'
-#pragma acc routine(StaticMemFunc)
+// CHECK-NEXT: gang clause
+// CHECK-NEXT: ConstantExpr{{.*}}'int'
+// CHECK-NEXT: value: Int 1
+#pragma acc routine(StaticMemFunc) gang(dim:Lambda())
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'void ()'
-#pragma acc routine(Lambda)
+// CHECK-NEXT: gang clause
+// CHECK-NEXT: ConstantExpr{{.*}}'int'
+// CHECK-NEXT: value: Int 1
+#pragma acc routine(Lambda) worker
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const S::(lambda at
+// CHECK-NEXT: worker clause
};
-#pragma acc routine(S::MemFunc)
+#pragma acc routine(S::MemFunc) gang(dim: 1)
// CHECK: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'void ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'S'
-#pragma acc routine(S::StaticMemFunc)
+// CHECK-NEXT: gang clause
+// CHECK-NEXT: ConstantExpr{{.*}}'int'
+// CHECK-NEXT: value: Int 1
+#pragma acc routine(S::StaticMemFunc) worker
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'void ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'S'
-#pragma acc routine(S::Lambda)
+// CHECK-NEXT: worker clause
+#pragma acc routine(S::Lambda) vector
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const S::(lambda at
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'S'
+// CHECK-NEXT: vector
template<typename T>
struct DepS {
T MemFunc();
static T StaticMemFunc();
- constexpr static auto Lambda = [](){};
+ constexpr static auto Lambda = [](){return 1;};
-#pragma acc routine(Lambda)
+#pragma acc routine(Lambda) gang(dim: Lambda())
// CHECK: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const auto'
-#pragma acc routine(MemFunc)
+// CHECK-NEXT: gang clause
+// CHECK-NEXT: CallExpr{{.*}}'<dependent type>'
+#pragma acc routine(MemFunc) worker
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'T ()'
-#pragma acc routine(StaticMemFunc)
+// CHECK-NEXT: worker clause
+#pragma acc routine(StaticMemFunc) seq
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'T ()'
+// CHECK-NEXT: seq clause
-#pragma acc routine(DepS::Lambda)
+#pragma acc routine(DepS::Lambda) gang(dim:1)
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const auto'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<T>'
-#pragma acc routine(DepS::MemFunc)
+// CHECK-NEXT: gang clause
+// CHECK-NEXT: ConstantExpr{{.*}}'int'
+// CHECK-NEXT: value: Int 1
+#pragma acc routine(DepS::MemFunc) gang
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'T ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<T>'
-#pragma acc routine(DepS::StaticMemFunc)
+// CHECK-NEXT: gang clause
+#pragma acc routine(DepS::StaticMemFunc) worker
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'T ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<T>'
+// CHECK-NEXT: worker clause
-#pragma acc routine(DepS<T>::Lambda)
+#pragma acc routine(DepS<T>::Lambda) vector
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const auto'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<T>'
-#pragma acc routine(DepS<T>::MemFunc)
+// CHECK-NEXT: vector clause
+#pragma acc routine(DepS<T>::MemFunc) seq
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'T ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<T>'
-#pragma acc routine(DepS<T>::StaticMemFunc)
+// CHECK-NEXT: seq clause
+#pragma acc routine(DepS<T>::StaticMemFunc) worker
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'T ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<T>'
+// CHECK-NEXT: worker clause
// Instantiation:
// CHECK: ClassTemplateSpecializationDecl{{.*}}struct DepS
@@ -118,68 +155,90 @@ struct DepS {
// CHECK: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const DepS<int>::
+// CHECK-NEXT: gang clause
+// CHECK-NEXT: ConstantExpr{{.*}}'int'
+// CHECK-NEXT: value: Int 1
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'int ()'
+// CHECK-NEXT: worker clause
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'int ()'
+// CHECK-NEXT: seq clause
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const DepS<int>::
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>'
+// CHECK-NEXT: gang clause
+// CHECK-NEXT: ConstantExpr{{.*}}'int'
+// CHECK-NEXT: value: Int 1
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'int ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>'
+// CHECK-NEXT: gang clause
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'int ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>'
+// CHECK-NEXT: worker clause
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const DepS<int>::
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>'
+// CHECK-NEXT: vector clause
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'int ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>'
+// CHECK-NEXT: seq clause
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'int ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>'
+// CHECK-NEXT: worker clause
};
-#pragma acc routine(DepS<int>::Lambda)
+#pragma acc routine(DepS<int>::Lambda) gang(dim:1)
// CHECK: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const DepS<int>::
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>'
-#pragma acc routine(DepS<int>::MemFunc)
+// CHECK-NEXT: gang clause
+// CHECK-NEXT: ConstantExpr{{.*}}'int'
+// CHECK-NEXT: value: Int 1
+#pragma acc routine(DepS<int>::MemFunc) worker
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'int ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>'
-#pragma acc routine(DepS<int>::StaticMemFunc)
+// CHECK-NEXT: worker clause
+#pragma acc routine(DepS<int>::StaticMemFunc) vector
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'int ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>'
+// CHECK-NEXT: vector clause
template<typename T>
void TemplFunc() {
-#pragma acc routine(T::MemFunc)
+#pragma acc routine(T::MemFunc) gang(dim:T::Lambda())
// CHECK: DeclStmt
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DependentScopeDeclRefExpr{{.*}}'<dependent type>'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'T'
-#pragma acc routine(T::StaticMemFunc)
+// CHECK-NEXT: gang clause
+// CHECK-NEXT: CallExpr{{.*}}'<dependent type>'
+#pragma acc routine(T::StaticMemFunc) worker
// CHECK-NEXT: DeclStmt
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DependentScopeDeclRefExpr{{.*}}'<dependent type>'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'T'
-#pragma acc routine(T::Lambda)
+// CHECK-NEXT: worker clause
+#pragma acc routine(T::Lambda) seq
// CHECK-NEXT: DeclStmt
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DependentScopeDeclRefExpr{{.*}}'<dependent type>'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'T'
+// CHECK-NEXT: seq clause
// Instantiation:
// CHECK: FunctionDecl{{.*}} TemplFunc 'void ()' implicit_instantiation
@@ -187,14 +246,21 @@ void TemplFunc() {
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'void ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'S'
+// CHECK-NEXT: gang clause
+// CHECK-NEXT: ConstantExpr{{.*}}'int'
+// CHECK-NEXT: value: Int 1
+
// CHECK-NEXT: DeclStmt
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'void ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'S'
+// CHECK-NEXT: worker clause
+
// CHECK-NEXT: DeclStmt
// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const S::(lambda at
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'S'
+// CHECK-NEXT: seq clause
}
void usage() {
@@ -202,3 +268,4 @@ void usage() {
TemplFunc<S>();
}
#endif
+
diff --git a/clang/test/SemaOpenACC/routine-construct-clauses.cpp b/clang/test/SemaOpenACC/routine-construct-clauses.cpp
new file mode 100644
index 0000000000000..87e566bfe81ce
--- /dev/null
+++ b/clang/test/SemaOpenACC/routine-construct-clauses.cpp
@@ -0,0 +1,119 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+void Func();
+
+#pragma acc routine(Func) worker
+#pragma acc routine(Func) vector
+#pragma acc routine(Func) seq
+#pragma acc routine(Func) gang
+
+// Only 1 of worker, vector, seq, gang.
+// expected-error at +2{{OpenACC clause 'vector' may not appear on the same construct as a 'worker' clause on a 'routine' construct}}
+// expected-note at +1{{previous clause is here}}
+#pragma acc routine(Func) worker vector
+// expected-error at +2{{OpenACC clause 'seq' may not appear on the same construct as a 'worker' clause on a 'routine' construct}}
+// expected-note at +1{{previous clause is here}}
+#pragma acc routine(Func) worker seq
+// expected-error at +2{{OpenACC clause 'gang' may not appear on the same construct as a 'worker' clause on a 'routine' construct}}
+// expected-note at +1{{previous clause is here}}
+#pragma acc routine(Func) worker gang
+// expected-error at +2{{OpenACC clause 'worker' may not appear on the same construct as a 'worker' clause on a 'routine' construct}}
+// expected-note at +1{{previous clause is here}}
+#pragma acc routine(Func) worker worker
+// expected-error at +2{{OpenACC clause 'worker' may not appear on the same construct as a 'vector' clause on a 'routine' construct}}
+// expected-note at +1{{previous clause is here}}
+#pragma acc routine(Func) vector worker
+// expected-error at +2{{OpenACC clause 'seq' may not appear on the same construct as a 'vector' clause on a 'routine' construct}}
+// expected-note at +1{{previous clause is here}}
+#pragma acc routine(Func) vector seq
+// expected-error at +2{{OpenACC clause 'gang' may not appear on the same construct as a 'vector' clause on a 'routine' construct}}
+// expected-note at +1{{previous clause is here}}
+#pragma acc routine(Func) vector gang
+// expected-error at +2{{OpenACC clause 'vector' may not appear on the same construct as a 'vector' clause on a 'routine' construct}}
+// expected-note at +1{{previous clause is here}}
+#pragma acc routine(Func) vector vector
+// expected-error at +2{{OpenACC clause 'worker' may not appear on the same construct as a 'seq' clause on a 'routine' construct}}
+// expected-note at +1{{previous clause is here}}
+#pragma acc routine(Func) seq worker
+// expected-error at +2{{OpenACC clause 'vector' may not appear on the same construct as a 'seq' clause on a 'routine' construct}}
+// expected-note at +1{{previous clause is here}}
+#pragma acc routine(Func) seq vector
+// expected-error at +2{{OpenACC clause 'gang' may not appear on the same construct as a 'seq' clause on a 'routine' construct}}
+// expected-note at +1{{previous clause is here}}
+#pragma acc routine(Func) seq gang
+// expected-error at +2{{OpenACC clause 'seq' may not appear on the same construct as a 'seq' clause on a 'routine' construct}}
+// expected-note at +1{{previous clause is here}}
+#pragma acc routine(Func) seq seq
+// expected-error at +2{{OpenACC clause 'worker' may not appear on the same construct as a 'gang' clause on a 'routine' construct}}
+// expected-note at +1{{previous clause is here}}
+#pragma acc routine(Func) gang worker
+// expected-error at +2{{OpenACC clause 'vector' may not appear on the same construct as a 'gang' clause on a 'routine' construct}}
+// expected-note at +1{{previous clause is here}}
+#pragma acc routine(Func) gang vector
+// expected-error at +2{{OpenACC clause 'seq' may not appear on the same construct as a 'gang' clause on a 'routine' construct}}
+// expected-note at +1{{previous clause is here}}
+#pragma acc routine(Func) gang seq
+// expected-error at +2{{OpenACC clause 'gang' may not appear on the same construct as a 'gang' clause on a 'routine' construct}}
+// expected-note at +1{{previous clause is here}}
+#pragma acc routine(Func) gang gang
+
+// only the 'dim' syntax for gang is legal.
+#pragma acc routine(Func) gang(dim:1)
+// expected-error at +1{{'num' argument on 'gang' clause is not permitted on a 'routine' construct}}
+#pragma acc routine(Func) gang(1)
+// expected-error at +1{{'num' argument on 'gang' clause is not permitted on a 'routine' construct}}
+#pragma acc routine(Func) gang(num:1)
+// expected-error at +1{{'static' argument on 'gang' clause is not permitted on a 'routine' construct}}
+#pragma acc routine(Func) gang(static:1)
+// expected-error at +2{{OpenACC 'gang' clause may have at most one 'dim' argument}}
+// expected-note at +1{{previous expression is here}}
+#pragma acc routine(Func) gang(dim:1, dim:2)
+
+// worker, vector, seq don't allow arguments.
+// expected-error at +1{{'num' argument on 'worker' clause is not permitted on a 'routine' construct}}
+#pragma acc routine(Func) worker(1)
+// expected-error at +1{{'length' argument on 'vector' clause is not permitted on a 'routine' construct}}
+#pragma acc routine(Func) vector(1)
+// expected-error at +1{{expected identifier}}
+#pragma acc routine(Func) seq(1)
+
+int getSomeInt();
+// dim must be a constant positive integer.
+// expected-error at +1{{argument to 'gang' clause dimension must be a constant expression}}
+#pragma acc routine(Func) gang(dim:getSomeInt())
+
+struct HasFuncs {
+static constexpr int Neg() { return -5; }
+static constexpr int Zero() { return 0; }
+static constexpr int One() { return 1; }
+static constexpr int Two() { return 2; }
+static constexpr int Three() { return 3; }
+static constexpr int Four() { return 4; }
+};
+// 'dim' must be 1, 2, or 3.
+// expected-error at +1{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to -5}}
+#pragma acc routine(Func) gang(dim:HasFuncs::Neg())
+// expected-error at +1{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 0}}
+#pragma acc routine(Func) gang(dim:HasFuncs::Zero())
+#pragma acc routine(Func) gang(dim:HasFuncs::One())
+#pragma acc routine(Func) gang(dim:HasFuncs::Two())
+#pragma acc routine(Func) gang(dim:HasFuncs::Three())
+// expected-error at +1{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 4}}
+#pragma acc routine(Func) gang(dim:HasFuncs::Four())
+
+template<typename T>
+struct DependentT {
+// expected-error at +1{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to -5}}
+#pragma acc routine(Func) gang(dim:T::Neg())
+// expected-error at +1{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 0}}
+#pragma acc routine(Func) gang(dim:T::Zero())
+#pragma acc routine(Func) gang(dim:T::One())
+#pragma acc routine(Func) gang(dim:T::Two())
+#pragma acc routine(Func) gang(dim:T::Three())
+// expected-error at +1{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 4}}
+#pragma acc routine(Func) gang(dim:T::Four())
+};
+
+void Inst() {
+ DependentT<HasFuncs> T;// expected-note{{in instantiation of}}
+}
diff --git a/clang/test/SemaOpenACC/routine-construct.cpp b/clang/test/SemaOpenACC/routine-construct.cpp
index f953a4537d0d6..3c7a31a2c0774 100644
--- a/clang/test/SemaOpenACC/routine-construct.cpp
+++ b/clang/test/SemaOpenACC/routine-construct.cpp
@@ -1,14 +1,14 @@
// RUN: %clang_cc1 %s -fopenacc -verify
// expected-error at +1{{use of undeclared identifier 'UnnamedYet'}}
-#pragma acc routine(UnnamedYet)
+#pragma acc routine(UnnamedYet) seq
void UnnamedYet();
// expected-error at +1{{use of undeclared identifier 'Invalid'}}
-#pragma acc routine(Invalid)
+#pragma acc routine(Invalid) seq
// Fine, since these are the same function.
void SameFunc();
-#pragma acc routine(SameFunc)
+#pragma acc routine(SameFunc) seq
void SameFunc();
void NoMagicStatic() {
@@ -16,12 +16,12 @@ void NoMagicStatic() {
}
// expected-error at -2{{function static variables are not permitted in functions to which an OpenACC 'routine' directive applies}}
// expected-note at +1{{'routine' construct is here}}
-#pragma acc routine(NoMagicStatic)
+#pragma acc routine(NoMagicStatic) seq
void NoMagicStatic2();
// expected-error at +4{{function static variables are not permitted in functions to which an OpenACC 'routine' directive applies}}
// expected-note at +1{{'routine' construct is here}}
-#pragma acc routine(NoMagicStatic2)
+#pragma acc routine(NoMagicStatic2) seq
void NoMagicStatic2() {
static int F = 1;
}
@@ -32,45 +32,45 @@ void HasMagicStaticLambda() {
};
// expected-error at -2{{function static variables are not permitted in functions to which an OpenACC 'routine' directive applies}}
// expected-note at +1{{'routine' construct is here}}
-#pragma acc routine (MSLambda)
+#pragma acc routine (MSLambda) seq
}
auto Lambda = [](){};
-#pragma acc routine(Lambda)
+#pragma acc routine(Lambda) seq
auto GenLambda = [](auto){};
// expected-error at +1{{OpenACC routine name 'GenLambda' names a set of overloads}}
-#pragma acc routine(GenLambda)
+#pragma acc routine(GenLambda) seq
// Variable?
int Variable;
// Plain function
int function();
-#pragma acc routine (function)
+#pragma acc routine (function) seq
// expected-error at +1{{OpenACC routine name 'Variable' does not name a function}}
-#pragma acc routine (Variable)
+#pragma acc routine (Variable) seq
// Var template?
template<typename T>
T VarTempl = 0;
// expected-error at +2{{use of variable template 'VarTempl' requires template arguments}}
// expected-note at -2{{template is declared here}}
-#pragma acc routine (VarTempl)
+#pragma acc routine (VarTempl) seq
// expected-error at +1{{OpenACC routine name 'VarTempl<int>' does not name a function}}
-#pragma acc routine (VarTempl<int>)
+#pragma acc routine (VarTempl<int>) seq
// Function in NS
namespace NS {
int NSFunc();
auto Lambda = [](){};
}
-#pragma acc routine(NS::NSFunc)
-#pragma acc routine(NS::Lambda)
+#pragma acc routine(NS::NSFunc) seq
+#pragma acc routine(NS::Lambda) seq
// Ambiguous Function
int ambig_func();
int ambig_func(int);
// expected-error at +1{{OpenACC routine name 'ambig_func' names a set of overloads}}
-#pragma acc routine (ambig_func)
+#pragma acc routine (ambig_func) seq
// Ambiguous in NS
namespace NS {
@@ -78,18 +78,18 @@ int ambig_func();
int ambig_func(int);
}
// expected-error at +1{{OpenACC routine name 'NS::ambig_func' names a set of overloads}}
-#pragma acc routine (NS::ambig_func)
+#pragma acc routine (NS::ambig_func) seq
// function template
template<typename T, typename U>
void templ_func();
// expected-error at +1{{OpenACC routine name 'templ_func' names a set of overloads}}
-#pragma acc routine(templ_func)
+#pragma acc routine(templ_func) seq
// expected-error at +1{{OpenACC routine name 'templ_func<int>' names a set of overloads}}
-#pragma acc routine(templ_func<int>)
+#pragma acc routine(templ_func<int>) seq
// expected-error at +1{{OpenACC routine name 'templ_func<int, float>' names a set of overloads}}
-#pragma acc routine(templ_func<int, float>)
+#pragma acc routine(templ_func<int, float>) seq
struct S {
void MemFunc();
@@ -110,46 +110,46 @@ struct S {
constexpr static auto Lambda = [](){};
-#pragma acc routine(S::MemFunc)
-#pragma acc routine(S::StaticMemFunc)
-#pragma acc routine(S::Lambda)
+#pragma acc routine(S::MemFunc) seq
+#pragma acc routine(S::StaticMemFunc) seq
+#pragma acc routine(S::Lambda) seq
// expected-error at +1{{OpenACC routine name 'S::TemplMemFunc' names a set of overloads}}
-#pragma acc routine(S::TemplMemFunc)
+#pragma acc routine(S::TemplMemFunc) seq
// expected-error at +1{{OpenACC routine name 'S::TemplStaticMemFunc' names a set of overloads}}
-#pragma acc routine(S::TemplStaticMemFunc)
+#pragma acc routine(S::TemplStaticMemFunc) seq
// expected-error at +1{{OpenACC routine name 'S::template TemplMemFunc<int>' names a set of overloads}}
-#pragma acc routine(S::template TemplMemFunc<int>)
+#pragma acc routine(S::template TemplMemFunc<int>) seq
// expected-error at +1{{OpenACC routine name 'S::template TemplStaticMemFunc<int>' names a set of overloads}}
-#pragma acc routine(S::template TemplStaticMemFunc<int>)
+#pragma acc routine(S::template TemplStaticMemFunc<int>) seq
// expected-error at +1{{OpenACC routine name 'S::MemFuncAmbig' names a set of overloads}}
-#pragma acc routine(S::MemFuncAmbig)
+#pragma acc routine(S::MemFuncAmbig) seq
// expected-error at +1{{OpenACC routine name 'S::TemplMemFuncAmbig' names a set of overloads}}
-#pragma acc routine(S::TemplMemFuncAmbig)
+#pragma acc routine(S::TemplMemFuncAmbig) seq
// expected-error at +1{{OpenACC routine name 'S::template TemplMemFuncAmbig<int>' names a set of overloads}}
-#pragma acc routine(S::template TemplMemFuncAmbig<int>)
+#pragma acc routine(S::template TemplMemFuncAmbig<int>) seq
// expected-error at +1{{OpenACC routine name 'S::Field' does not name a function}}
-#pragma acc routine(S::Field)
+#pragma acc routine(S::Field) seq
};
-#pragma acc routine(S::MemFunc)
-#pragma acc routine(S::StaticMemFunc)
-#pragma acc routine(S::Lambda)
+#pragma acc routine(S::MemFunc) seq
+#pragma acc routine(S::StaticMemFunc) seq
+#pragma acc routine(S::Lambda) seq
// expected-error at +1{{OpenACC routine name 'S::TemplMemFunc' names a set of overloads}}
-#pragma acc routine(S::TemplMemFunc)
+#pragma acc routine(S::TemplMemFunc) seq
// expected-error at +1{{OpenACC routine name 'S::TemplStaticMemFunc' names a set of overloads}}
-#pragma acc routine(S::TemplStaticMemFunc)
+#pragma acc routine(S::TemplStaticMemFunc) seq
// expected-error at +1{{OpenACC routine name 'S::template TemplMemFunc<int>' names a set of overloads}}
-#pragma acc routine(S::template TemplMemFunc<int>)
+#pragma acc routine(S::template TemplMemFunc<int>) seq
// expected-error at +1{{OpenACC routine name 'S::template TemplStaticMemFunc<int>' names a set of overloads}}
-#pragma acc routine(S::template TemplStaticMemFunc<int>)
+#pragma acc routine(S::template TemplStaticMemFunc<int>) seq
// expected-error at +1{{OpenACC routine name 'S::MemFuncAmbig' names a set of overloads}}
-#pragma acc routine(S::MemFuncAmbig)
+#pragma acc routine(S::MemFuncAmbig) seq
// expected-error at +1{{OpenACC routine name 'S::TemplMemFuncAmbig' names a set of overloads}}
-#pragma acc routine(S::TemplMemFuncAmbig)
+#pragma acc routine(S::TemplMemFuncAmbig) seq
// expected-error at +1{{OpenACC routine name 'S::template TemplMemFuncAmbig<int>' names a set of overloads}}
-#pragma acc routine(S::template TemplMemFuncAmbig<int>)
+#pragma acc routine(S::template TemplMemFuncAmbig<int>) seq
// expected-error at +1{{OpenACC routine name 'S::Field' does not name a function}}
-#pragma acc routine(S::Field)
+#pragma acc routine(S::Field) seq
template<typename T>
struct DepS { // #DEPS
@@ -172,49 +172,49 @@ struct DepS { // #DEPS
constexpr static auto Lambda = [](){};
// expected-error at +2{{non-const static data member must be initialized out of line}}
// expected-note@#DEPSInst{{in instantiation of template class}}
- static auto LambdaBroken = [](){};
+ static auto LambdaBroken = [](){};
-#pragma acc routine(DepS::MemFunc)
-#pragma acc routine(DepS::StaticMemFunc)
-#pragma acc routine(DepS::Lambda)
-#pragma acc routine(DepS::LambdaBroken)
+#pragma acc routine(DepS::MemFunc) seq
+#pragma acc routine(DepS::StaticMemFunc) seq
+#pragma acc routine(DepS::Lambda) seq
+#pragma acc routine(DepS::LambdaBroken) seq
// expected-error at +1{{OpenACC routine name 'DepS<T>::TemplMemFunc' names a set of overloads}}
-#pragma acc routine(DepS::TemplMemFunc)
+#pragma acc routine(DepS::TemplMemFunc) seq
// expected-error at +1{{OpenACC routine name 'DepS<T>::TemplStaticMemFunc' names a set of overloads}}
-#pragma acc routine(DepS::TemplStaticMemFunc)
+#pragma acc routine(DepS::TemplStaticMemFunc) seq
// expected-error at +1{{OpenACC routine name 'DepS<T>::template TemplMemFunc<int>' names a set of overloads}}
-#pragma acc routine(DepS::template TemplMemFunc<int>)
+#pragma acc routine(DepS::template TemplMemFunc<int>) seq
// expected-error at +1{{OpenACC routine name 'DepS<T>::template TemplStaticMemFunc<int>' names a set of overloads}}
-#pragma acc routine(DepS::template TemplStaticMemFunc<int>)
+#pragma acc routine(DepS::template TemplStaticMemFunc<int>) seq
// expected-error at +1{{OpenACC routine name 'DepS<T>::MemFuncAmbig' names a set of overloads}}
-#pragma acc routine(DepS::MemFuncAmbig)
+#pragma acc routine(DepS::MemFuncAmbig) seq
// expected-error at +1{{OpenACC routine name 'DepS<T>::TemplMemFuncAmbig' names a set of overloads}}
-#pragma acc routine(DepS::TemplMemFuncAmbig)
+#pragma acc routine(DepS::TemplMemFuncAmbig) seq
// expected-error at +1{{OpenACC routine name 'DepS<T>::template TemplMemFuncAmbig<int>' names a set of overloads}}
-#pragma acc routine(DepS::template TemplMemFuncAmbig<int>)
+#pragma acc routine(DepS::template TemplMemFuncAmbig<int>) seq
// expected-error at +1{{OpenACC routine name 'DepS<T>::Field' does not name a function}}
-#pragma acc routine(DepS::Field)
+#pragma acc routine(DepS::Field) seq
-#pragma acc routine(DepS<T>::MemFunc)
-#pragma acc routine(DepS<T>::StaticMemFunc)
-#pragma acc routine(DepS<T>::Lambda)
-#pragma acc routine(DepS<T>::LambdaBroken)
+#pragma acc routine(DepS<T>::MemFunc) seq
+#pragma acc routine(DepS<T>::StaticMemFunc) seq
+#pragma acc routine(DepS<T>::Lambda) seq
+#pragma acc routine(DepS<T>::LambdaBroken) seq
// expected-error at +1{{OpenACC routine name 'DepS<T>::TemplMemFunc' names a set of overloads}}
-#pragma acc routine(DepS<T>::TemplMemFunc)
+#pragma acc routine(DepS<T>::TemplMemFunc) seq
// expected-error at +1{{OpenACC routine name 'DepS<T>::TemplStaticMemFunc' names a set of overloads}}
-#pragma acc routine(DepS<T>::TemplStaticMemFunc)
+#pragma acc routine(DepS<T>::TemplStaticMemFunc) seq
// expected-error at +1{{OpenACC routine name 'DepS<T>::template TemplMemFunc<int>' names a set of overloads}}
-#pragma acc routine(DepS<T>::template TemplMemFunc<int>)
+#pragma acc routine(DepS<T>::template TemplMemFunc<int>) seq
// expected-error at +1{{OpenACC routine name 'DepS<T>::template TemplStaticMemFunc<int>' names a set of overloads}}
-#pragma acc routine(DepS<T>::template TemplStaticMemFunc<int>)
+#pragma acc routine(DepS<T>::template TemplStaticMemFunc<int>) seq
// expected-error at +1{{OpenACC routine name 'DepS<T>::MemFuncAmbig' names a set of overloads}}
-#pragma acc routine(DepS<T>::MemFuncAmbig)
+#pragma acc routine(DepS<T>::MemFuncAmbig) seq
// expected-error at +1{{OpenACC routine name 'DepS<T>::TemplMemFuncAmbig' names a set of overloads}}
-#pragma acc routine(DepS<T>::TemplMemFuncAmbig)
+#pragma acc routine(DepS<T>::TemplMemFuncAmbig) seq
// expected-error at +1{{OpenACC routine name 'DepS<T>::template TemplMemFuncAmbig<int>' names a set of overloads}}
-#pragma acc routine(DepS<T>::template TemplMemFuncAmbig<int>)
+#pragma acc routine(DepS<T>::template TemplMemFuncAmbig<int>) seq
// expected-error at +1{{OpenACC routine name 'DepS<T>::Field' does not name a function}}
-#pragma acc routine(DepS<T>::Field)
+#pragma acc routine(DepS<T>::Field) seq
};
void Inst() {
@@ -224,170 +224,170 @@ void Inst() {
//expected-error at +2{{use of class template 'DepS' requires template arguments}}
// expected-note@#DEPS{{template is declared here}}
-#pragma acc routine(DepS::Lambda)
+#pragma acc routine(DepS::Lambda) seq
//expected-error at +2{{use of class template 'DepS' requires template arguments}}
// expected-note@#DEPS{{template is declared here}}
-#pragma acc routine(DepS::MemFunc)
+#pragma acc routine(DepS::MemFunc) seq
//expected-error at +2{{use of class template 'DepS' requires template arguments}}
// expected-note@#DEPS{{template is declared here}}
-#pragma acc routine(DepS::StaticMemFunc)
+#pragma acc routine(DepS::StaticMemFunc) seq
//expected-error at +2{{use of class template 'DepS' requires template arguments}}
// expected-note@#DEPS{{template is declared here}}
-#pragma acc routine(DepS::TemplMemFunc)
+#pragma acc routine(DepS::TemplMemFunc) seq
//expected-error at +2{{use of class template 'DepS' requires template arguments}}
// expected-note@#DEPS{{template is declared here}}
-#pragma acc routine(DepS::TemplStaticMemFunc)
+#pragma acc routine(DepS::TemplStaticMemFunc) seq
//expected-error at +2{{use of class template 'DepS' requires template arguments}}
// expected-note@#DEPS{{template is declared here}}
-#pragma acc routine(DepS::TemplMemFunc<int>)
+#pragma acc routine(DepS::TemplMemFunc<int>) seq
//expected-error at +2{{use of class template 'DepS' requires template arguments}}
// expected-note@#DEPS{{template is declared here}}
-#pragma acc routine(DepS::TemplStaticMemFunc<int>)
+#pragma acc routine(DepS::TemplStaticMemFunc<int>) seq
//expected-error at +2{{use of class template 'DepS' requires template arguments}}
// expected-note@#DEPS{{template is declared here}}
-#pragma acc routine(DepS::MemFuncAmbig)
+#pragma acc routine(DepS::MemFuncAmbig) seq
//expected-error at +2{{use of class template 'DepS' requires template arguments}}
// expected-note@#DEPS{{template is declared here}}
-#pragma acc routine(DepS::TemplMemFuncAmbig)
+#pragma acc routine(DepS::TemplMemFuncAmbig) seq
//expected-error at +2{{use of class template 'DepS' requires template arguments}}
// expected-note@#DEPS{{template is declared here}}
-#pragma acc routine(DepS::TemplMemFuncAmbig<int>)
+#pragma acc routine(DepS::TemplMemFuncAmbig<int>) seq
//expected-error at +2{{use of class template 'DepS' requires template arguments}}
// expected-note@#DEPS{{template is declared here}}
-#pragma acc routine(DepS::Field)
+#pragma acc routine(DepS::Field) seq
//expected-error at +1{{use of undeclared identifier 'T'}}
-#pragma acc routine(DepS<T>::Lambda)
+#pragma acc routine(DepS<T>::Lambda) seq
//expected-error at +1{{use of undeclared identifier 'T'}}
-#pragma acc routine(DepS<T>::MemFunc)
+#pragma acc routine(DepS<T>::MemFunc) seq
//expected-error at +1{{use of undeclared identifier 'T'}}
-#pragma acc routine(DepS<T>::StaticMemFunc)
+#pragma acc routine(DepS<T>::StaticMemFunc) seq
//expected-error at +1{{use of undeclared identifier 'T'}}
-#pragma acc routine(DepS<T>::TemplMemFunc)
+#pragma acc routine(DepS<T>::TemplMemFunc) seq
//expected-error at +1{{use of undeclared identifier 'T'}}
-#pragma acc routine(DepS<T>::TemplStaticMemFunc)
+#pragma acc routine(DepS<T>::TemplStaticMemFunc) seq
//expected-error at +1{{use of undeclared identifier 'T'}}
-#pragma acc routine(DepS<T>::TemplMemFunc<int>)
+#pragma acc routine(DepS<T>::TemplMemFunc<int>) seq
//expected-error at +1{{use of undeclared identifier 'T'}}
-#pragma acc routine(DepS<T>::TemplStaticMemFunc<int>)
+#pragma acc routine(DepS<T>::TemplStaticMemFunc<int>) seq
//expected-error at +1{{use of undeclared identifier 'T'}}
-#pragma acc routine(DepS<T>::MemFuncAmbig)
+#pragma acc routine(DepS<T>::MemFuncAmbig) seq
//expected-error at +1{{use of undeclared identifier 'T'}}
-#pragma acc routine(DepS<T>::TemplMemFuncAmbig)
+#pragma acc routine(DepS<T>::TemplMemFuncAmbig) seq
//expected-error at +1{{use of undeclared identifier 'T'}}
-#pragma acc routine(DepS<T>::TemplMemFuncAmbig<int>)
+#pragma acc routine(DepS<T>::TemplMemFuncAmbig<int>) seq
//expected-error at +1{{use of undeclared identifier 'T'}}
-#pragma acc routine(DepS<T>::Field)
+#pragma acc routine(DepS<T>::Field) seq
-#pragma acc routine(DepS<int>::Lambda)
-#pragma acc routine(DepS<int>::LambdaBroken)
-#pragma acc routine(DepS<int>::MemFunc)
-#pragma acc routine(DepS<int>::StaticMemFunc)
+#pragma acc routine(DepS<int>::Lambda) seq
+#pragma acc routine(DepS<int>::LambdaBroken) seq
+#pragma acc routine(DepS<int>::MemFunc) seq
+#pragma acc routine(DepS<int>::StaticMemFunc) seq
// expected-error at +1{{OpenACC routine name 'DepS<int>::TemplMemFunc' names a set of overloads}}
-#pragma acc routine(DepS<int>::TemplMemFunc)
+#pragma acc routine(DepS<int>::TemplMemFunc) seq
// expected-error at +1{{OpenACC routine name 'DepS<int>::TemplStaticMemFunc' names a set of overloads}}
-#pragma acc routine(DepS<int>::TemplStaticMemFunc)
+#pragma acc routine(DepS<int>::TemplStaticMemFunc) seq
// expected-error at +1{{OpenACC routine name 'DepS<int>::TemplMemFunc<int>' names a set of overloads}}
-#pragma acc routine(DepS<int>::TemplMemFunc<int>)
+#pragma acc routine(DepS<int>::TemplMemFunc<int>) seq
// expected-error at +1{{OpenACC routine name 'DepS<int>::TemplStaticMemFunc<int>' names a set of overloads}}
-#pragma acc routine(DepS<int>::TemplStaticMemFunc<int>)
+#pragma acc routine(DepS<int>::TemplStaticMemFunc<int>) seq
// expected-error at +1{{OpenACC routine name 'DepS<int>::MemFuncAmbig' names a set of overloads}}
-#pragma acc routine(DepS<int>::MemFuncAmbig)
+#pragma acc routine(DepS<int>::MemFuncAmbig) seq
// expected-error at +1{{OpenACC routine name 'DepS<int>::TemplMemFuncAmbig' names a set of overloads}}
-#pragma acc routine(DepS<int>::TemplMemFuncAmbig)
+#pragma acc routine(DepS<int>::TemplMemFuncAmbig) seq
// expected-error at +1{{OpenACC routine name 'DepS<int>::TemplMemFuncAmbig<int>' names a set of overloads}}
-#pragma acc routine(DepS<int>::TemplMemFuncAmbig<int>)
+#pragma acc routine(DepS<int>::TemplMemFuncAmbig<int>) seq
// expected-error at +1{{OpenACC routine name 'DepS<int>::Field' does not name a function}}
-#pragma acc routine(DepS<int>::Field)
+#pragma acc routine(DepS<int>::Field) seq
template<typename T>
void TemplFunc() {
-#pragma acc routine(T::MemFunc)
-#pragma acc routine(T::StaticMemFunc)
-#pragma acc routine(T::Lambda)
+#pragma acc routine(T::MemFunc) seq
+#pragma acc routine(T::StaticMemFunc) seq
+#pragma acc routine(T::Lambda) seq
// expected-error at +1{{OpenACC routine name 'S::TemplMemFunc' names a set of overloads}}
-#pragma acc routine(T::TemplMemFunc)
+#pragma acc routine(T::TemplMemFunc) seq
// expected-error at +1{{OpenACC routine name 'S::TemplStaticMemFunc' names a set of overloads}}
-#pragma acc routine(T::TemplStaticMemFunc)
+#pragma acc routine(T::TemplStaticMemFunc) seq
// expected-error at +1{{OpenACC routine name 'S::template TemplMemFunc<int>' names a set of overloads}}
-#pragma acc routine(T::template TemplMemFunc<int>)
+#pragma acc routine(T::template TemplMemFunc<int>) seq
// expected-error at +1{{OpenACC routine name 'S::template TemplStaticMemFunc<int>' names a set of overloads}}
-#pragma acc routine(T::template TemplStaticMemFunc<int>)
+#pragma acc routine(T::template TemplStaticMemFunc<int>) seq
// expected-error at +1{{OpenACC routine name 'S::MemFuncAmbig' names a set of overloads}}
-#pragma acc routine(T::MemFuncAmbig)
+#pragma acc routine(T::MemFuncAmbig) seq
// expected-error at +1{{OpenACC routine name 'S::TemplMemFuncAmbig' names a set of overloads}}
-#pragma acc routine(T::TemplMemFuncAmbig)
+#pragma acc routine(T::TemplMemFuncAmbig) seq
// expected-error at +1{{OpenACC routine name 'S::template TemplMemFuncAmbig<int>' names a set of overloads}}
-#pragma acc routine(T::template TemplMemFuncAmbig<int>)
+#pragma acc routine(T::template TemplMemFuncAmbig<int>) seq
// expected-error at +1{{OpenACC routine name 'S::Field' does not name a function}}
-#pragma acc routine(T::Field)
+#pragma acc routine(T::Field) seq
}
template <typename T>
struct DepRefersToT {
-#pragma acc routine(T::MemFunc)
-#pragma acc routine(T::StaticMemFunc)
-#pragma acc routine(T::Lambda)
+#pragma acc routine(T::MemFunc) seq
+#pragma acc routine(T::StaticMemFunc) seq
+#pragma acc routine(T::Lambda) seq
// expected-error at +1{{OpenACC routine name 'S::TemplMemFunc' names a set of overloads}}
-#pragma acc routine(T::TemplMemFunc)
+#pragma acc routine(T::TemplMemFunc) seq
// expected-error at +1{{OpenACC routine name 'S::TemplStaticMemFunc' names a set of overloads}}
-#pragma acc routine(T::TemplStaticMemFunc)
+#pragma acc routine(T::TemplStaticMemFunc) seq
// expected-error at +1{{OpenACC routine name 'S::template TemplMemFunc<int>' names a set of overloads}}
-#pragma acc routine(T::template TemplMemFunc<int>)
+#pragma acc routine(T::template TemplMemFunc<int>) seq
// expected-error at +1{{OpenACC routine name 'S::template TemplStaticMemFunc<int>' names a set of overloads}}
-#pragma acc routine(T::template TemplStaticMemFunc<int>)
+#pragma acc routine(T::template TemplStaticMemFunc<int>) seq
// expected-error at +1{{OpenACC routine name 'S::MemFuncAmbig' names a set of overloads}}
-#pragma acc routine(T::MemFuncAmbig)
+#pragma acc routine(T::MemFuncAmbig) seq
// expected-error at +1{{OpenACC routine name 'S::TemplMemFuncAmbig' names a set of overloads}}
-#pragma acc routine(T::TemplMemFuncAmbig)
+#pragma acc routine(T::TemplMemFuncAmbig) seq
// expected-error at +1{{OpenACC routine name 'S::template TemplMemFuncAmbig<int>' names a set of overloads}}
-#pragma acc routine(T::template TemplMemFuncAmbig<int>)
+#pragma acc routine(T::template TemplMemFuncAmbig<int>) seq
// expected-error at +1{{OpenACC routine name 'S::Field' does not name a function}}
-#pragma acc routine(T::Field)
+#pragma acc routine(T::Field) seq
void MemFunc() {
-#pragma acc routine(T::MemFunc)
-#pragma acc routine(T::StaticMemFunc)
-#pragma acc routine(T::Lambda)
+#pragma acc routine(T::MemFunc) seq
+#pragma acc routine(T::StaticMemFunc) seq
+#pragma acc routine(T::Lambda) seq
// expected-error at +1{{OpenACC routine name 'S::TemplMemFunc' names a set of overloads}}
-#pragma acc routine(T::TemplMemFunc)
+#pragma acc routine(T::TemplMemFunc) seq
// expected-error at +1{{OpenACC routine name 'S::TemplStaticMemFunc' names a set of overloads}}
-#pragma acc routine(T::TemplStaticMemFunc)
+#pragma acc routine(T::TemplStaticMemFunc) seq
// expected-error at +1{{OpenACC routine name 'S::template TemplMemFunc<int>' names a set of overloads}}
-#pragma acc routine(T::template TemplMemFunc<int>)
+#pragma acc routine(T::template TemplMemFunc<int>) seq
// expected-error at +1{{OpenACC routine name 'S::template TemplStaticMemFunc<int>' names a set of overloads}}
-#pragma acc routine(T::template TemplStaticMemFunc<int>)
+#pragma acc routine(T::template TemplStaticMemFunc<int>) seq
// expected-error at +1{{OpenACC routine name 'S::MemFuncAmbig' names a set of overloads}}
-#pragma acc routine(T::MemFuncAmbig)
+#pragma acc routine(T::MemFuncAmbig) seq
// expected-error at +1{{OpenACC routine name 'S::TemplMemFuncAmbig' names a set of overloads}}
-#pragma acc routine(T::TemplMemFuncAmbig)
+#pragma acc routine(T::TemplMemFuncAmbig) seq
// expected-error at +1{{OpenACC routine name 'S::template TemplMemFuncAmbig<int>' names a set of overloads}}
-#pragma acc routine(T::template TemplMemFuncAmbig<int>)
+#pragma acc routine(T::template TemplMemFuncAmbig<int>) seq
// expected-error at +1{{OpenACC routine name 'S::Field' does not name a function}}
-#pragma acc routine(T::Field)
+#pragma acc routine(T::Field) seq
}
template<typename U>
void TemplMemFunc() {
-#pragma acc routine(T::MemFunc)
-#pragma acc routine(T::StaticMemFunc)
-#pragma acc routine(T::Lambda)
+#pragma acc routine(T::MemFunc) seq
+#pragma acc routine(T::StaticMemFunc) seq
+#pragma acc routine(T::Lambda) seq
// expected-error at +1{{OpenACC routine name 'S::TemplMemFunc' names a set of overloads}}
-#pragma acc routine(T::TemplMemFunc)
+#pragma acc routine(T::TemplMemFunc) seq
// expected-error at +1{{OpenACC routine name 'S::TemplStaticMemFunc' names a set of overloads}}
-#pragma acc routine(T::TemplStaticMemFunc)
+#pragma acc routine(T::TemplStaticMemFunc) seq
// expected-error at +1{{OpenACC routine name 'S::template TemplMemFunc<int>' names a set of overloads}}
-#pragma acc routine(T::template TemplMemFunc<int>)
+#pragma acc routine(T::template TemplMemFunc<int>) seq
// expected-error at +1{{OpenACC routine name 'S::template TemplStaticMemFunc<int>' names a set of overloads}}
-#pragma acc routine(T::template TemplStaticMemFunc<int>)
+#pragma acc routine(T::template TemplStaticMemFunc<int>) seq
// expected-error at +1{{OpenACC routine name 'S::MemFuncAmbig' names a set of overloads}}
-#pragma acc routine(T::MemFuncAmbig)
+#pragma acc routine(T::MemFuncAmbig) seq
// expected-error at +1{{OpenACC routine name 'S::TemplMemFuncAmbig' names a set of overloads}}
-#pragma acc routine(T::TemplMemFuncAmbig)
+#pragma acc routine(T::TemplMemFuncAmbig) seq
// expected-error at +1{{OpenACC routine name 'S::template TemplMemFuncAmbig<int>' names a set of overloads}}
-#pragma acc routine(T::template TemplMemFuncAmbig<int>)
+#pragma acc routine(T::template TemplMemFuncAmbig<int>) seq
// expected-error at +1{{OpenACC routine name 'S::Field' does not name a function}}
-#pragma acc routine(T::Field)
+#pragma acc routine(T::Field) seq
}
};
diff --git a/clang/test/SemaOpenACC/unimplemented-construct.c b/clang/test/SemaOpenACC/unimplemented-construct.c
index ec11792c4f99b..4e29189de13b6 100644
--- a/clang/test/SemaOpenACC/unimplemented-construct.c
+++ b/clang/test/SemaOpenACC/unimplemented-construct.c
@@ -1,38 +1,38 @@
// RUN: %clang_cc1 %s -verify -fopenacc
// expected-warning at +1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}}
-#pragma acc routine
+#pragma acc routine seq
struct S {
// expected-warning at +1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}}
-#pragma acc routine
+#pragma acc routine seq
int foo;
};
void func() {
// expected-warning at +1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}}
-#pragma acc routine
+#pragma acc routine seq
int foo;
// expected-warning at +1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}}
-#pragma acc routine
+#pragma acc routine seq
{
// expected-warning at +1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}}
-#pragma acc routine
+#pragma acc routine seq
{
// expected-warning at +1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}}
-#pragma acc routine
+#pragma acc routine seq
}
}
// expected-warning at +1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}}
-#pragma acc routine
+#pragma acc routine seq
while(0){}
// expected-warning at +1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}}
-#pragma acc routine
+#pragma acc routine seq
for(;;){}
// expected-warning at +1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}}
-#pragma acc routine
+#pragma acc routine seq
};
More information about the cfe-commits
mailing list