[clang] 79079c9 - [OpenACC] Finish implementing 'routine' AST/Sema.
via cfe-commits
cfe-commits at lists.llvm.org
Fri Mar 21 08:57:59 PDT 2025
Author: erichkeane
Date: 2025-03-21T08:57:54-07:00
New Revision: 79079c94699a4210bd500556d2864cbdf4a37e78
URL: https://github.com/llvm/llvm-project/commit/79079c94699a4210bd500556d2864cbdf4a37e78
DIFF: https://github.com/llvm/llvm-project/commit/79079c94699a4210bd500556d2864cbdf4a37e78.diff
LOG: [OpenACC] Finish implementing 'routine' AST/Sema.
This is the last item of the OpenACC 3.3 spec. It includes the
implicit-name version of 'routine', plus significant refactorings to
make the two work together. The implicit name version is represented as
an attribute on the function call. This patch also implements the
clauses for the implicit-name version, as well as the A.3.4 warning.
Added:
Modified:
clang/include/clang/AST/Attr.h
clang/include/clang/AST/DeclOpenACC.h
clang/include/clang/AST/OpenACCClause.h
clang/include/clang/AST/TextNodeDumper.h
clang/include/clang/Basic/Attr.td
clang/include/clang/Basic/DiagnosticSemaKinds.td
clang/include/clang/Parse/Parser.h
clang/include/clang/Sema/SemaOpenACC.h
clang/include/clang/Serialization/ASTRecordReader.h
clang/include/clang/Serialization/ASTRecordWriter.h
clang/lib/AST/DeclOpenACC.cpp
clang/lib/AST/DeclPrinter.cpp
clang/lib/AST/OpenACCClause.cpp
clang/lib/AST/TextNodeDumper.cpp
clang/lib/Parse/ParseDecl.cpp
clang/lib/Parse/ParseDeclCXX.cpp
clang/lib/Parse/ParseOpenACC.cpp
clang/lib/Parse/Parser.cpp
clang/lib/Sema/SemaDecl.cpp
clang/lib/Sema/SemaOpenACC.cpp
clang/lib/Sema/SemaOpenACCClause.cpp
clang/lib/Sema/SemaTemplateInstantiate.cpp
clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
clang/lib/Serialization/ASTReader.cpp
clang/lib/Serialization/ASTWriter.cpp
clang/test/AST/ast-print-openacc-routine-construct.cpp
clang/test/ParserOpenACC/parse-clauses.c
clang/test/ParserOpenACC/parse-constructs.c
clang/test/ParserOpenACC/parse-constructs.cpp
clang/test/SemaOpenACC/routine-construct-ast.cpp
clang/test/SemaOpenACC/routine-construct-clauses.cpp
clang/test/SemaOpenACC/routine-construct.cpp
clang/utils/TableGen/ClangAttrEmitter.cpp
Removed:
clang/test/SemaOpenACC/unimplemented-construct.c
################################################################################
diff --git a/clang/include/clang/AST/Attr.h b/clang/include/clang/AST/Attr.h
index 3365ebe4d9012..994f236337b99 100644
--- a/clang/include/clang/AST/Attr.h
+++ b/clang/include/clang/AST/Attr.h
@@ -38,6 +38,7 @@ class ASTContext;
class AttributeCommonInfo;
class FunctionDecl;
class OMPTraitInfo;
+class OpenACCClause;
/// Attr - This represents one attribute.
class Attr : public AttributeCommonInfo {
diff --git a/clang/include/clang/AST/DeclOpenACC.h b/clang/include/clang/AST/DeclOpenACC.h
index 9e99061ffc47f..26cf721561fb1 100644
--- a/clang/include/clang/AST/DeclOpenACC.h
+++ b/clang/include/clang/AST/DeclOpenACC.h
@@ -102,6 +102,8 @@ class OpenACCDeclareDecl final
static bool classofKind(Kind K) { return K == OpenACCDeclare; }
};
+// Reprents a 'routine' directive with a name. When this has no name, it is
+// represented as an attribute.
class OpenACCRoutineDecl final
: public OpenACCConstructDecl,
private llvm::TrailingObjects<OpenACCRoutineDecl, const OpenACCClause *> {
@@ -129,6 +131,8 @@ class OpenACCRoutineDecl final
: OpenACCConstructDecl(OpenACCRoutine, DC, OpenACCDirectiveKind::Routine,
StartLoc, DirLoc, EndLoc),
FuncRef(FuncRef), ParensLoc(LParenLoc, RParenLoc) {
+ assert(LParenLoc.isValid() &&
+ "Cannot represent implicit name with this declaration");
// Initialize the trailing storage.
std::uninitialized_copy(Clauses.begin(), Clauses.end(),
getTrailingObjects<const OpenACCClause *>());
@@ -148,13 +152,10 @@ class OpenACCRoutineDecl final
static bool classofKind(Kind K) { return K == OpenACCRoutine; }
const Expr *getFunctionReference() const { return FuncRef; }
-
Expr *getFunctionReference() { return FuncRef; }
SourceLocation getLParenLoc() const { return ParensLoc.getBegin(); }
SourceLocation getRParenLoc() const { return ParensLoc.getEnd(); }
-
- bool hasNameSpecified() const { return !ParensLoc.getBegin().isInvalid(); }
};
} // namespace clang
diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h
index 049389229b8d5..4c5fe03a34361 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -251,6 +251,12 @@ class OpenACCBindClause final : public OpenACCClauseWithParams {
}
};
+bool operator==(const OpenACCBindClause &LHS, const OpenACCBindClause &RHS);
+inline bool operator!=(const OpenACCBindClause &LHS,
+ const OpenACCBindClause &RHS) {
+ return !(LHS == RHS);
+}
+
using DeviceTypeArgument = std::pair<IdentifierInfo *, SourceLocation>;
/// A 'device_type' or 'dtype' clause, takes a list of either an 'asterisk' or
/// an identifier. The 'asterisk' means 'the rest'.
diff --git a/clang/include/clang/AST/TextNodeDumper.h b/clang/include/clang/AST/TextNodeDumper.h
index ad2d4a7a973b3..ea3a0f058a8ed 100644
--- a/clang/include/clang/AST/TextNodeDumper.h
+++ b/clang/include/clang/AST/TextNodeDumper.h
@@ -426,6 +426,7 @@ class TextNodeDumper
void VisitOpenACCAsteriskSizeExpr(const OpenACCAsteriskSizeExpr *S);
void VisitOpenACCDeclareDecl(const OpenACCDeclareDecl *D);
void VisitOpenACCRoutineDecl(const OpenACCRoutineDecl *D);
+ void VisitOpenACCRoutineDeclAttr(const OpenACCRoutineDeclAttr *A);
void VisitEmbedExpr(const EmbedExpr *S);
void VisitAtomicExpr(const AtomicExpr *AE);
void VisitConvertVectorExpr(const ConvertVectorExpr *S);
diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index b8cd3475bb88a..1e9db55fbe301 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -693,6 +693,10 @@ class Attr {
// content. Eg) It parses 3 args, but semantically takes 4 args. Opts out of
// common attribute error checking.
bit HasCustomParsing = 0;
+ // Set to true if this attribute requires custom serialization after the
+ // typical attribute serialization. This will cause tablegen to emit a call to
+ // ASTRecordWriter::Add<Name>Attr and ASTRecordReader::read<Name>Attr.
+ bit HasCustomSerialization = 0;
// Set to true if all of the attribute's arguments should be parsed in an
// unevaluated context.
bit ParseArgumentsAsUnevaluated = 0;
@@ -4991,12 +4995,30 @@ def Atomic : StmtAttr {
}
def OpenACCRoutineAnnot : InheritableAttr {
- // This attribute is used to mark that a function is targetted by a `routine`
- // directive, so it dones't have a spelling and is always implicit.
+ // This attribute is used to mark that a function is targeted by a `routine`
+ // directive with a name for the purposes of checking the declaration later.
+ // We don't really need a link back to the declaration, as location is
+ // sufficient.
+ // We abuse the source locations on this a little, since we need two locations
+ // for various diagnostic purposes. The 'begin' location is the location of
+ // the Routine directive. We are using the 'end' location for any 'bind'
+ // clauses, since this is needed for a diagnostic.
let Spellings = [];
let Subjects = SubjectList<[Function]>;
let Documentation = [InternalOnly];
+}
+
+def OpenACCRoutineDecl :InheritableAttr {
+ // This attribute represents the 'routine' directive when spelled without a
+ // 'name'.
+ let Spellings = [Pragma<"acc", "routine">];
+ let Subjects = SubjectList<[Function]>;
+ let SemaHandler = 0;
+ let HasCustomParsing = 1;
+ let HasCustomSerialization = 1;
+ let Documentation = [InternalOnly];
let AdditionalMembers = [{
- SourceLocation BindClause;
+ llvm::SmallVector<const OpenACCClause *> Clauses;
+ void printPrettyPragma(raw_ostream &OS, const PrintingPolicy &Policy) const;
}];
}
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index eed06657536f4..a3becb7520a98 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -13040,6 +13040,16 @@ def err_acc_magic_static_in_routine
def err_acc_duplicate_bind
: Error<"multiple 'routine' directives with 'bind' clauses are not "
"permitted to refer to the same function">;
+def err_acc_duplicate_unnamed_bind
+ : Error<"OpenACC 'bind' clause on a declaration must bind to the same name "
+ "as previous bind clauses">;
+def warn_acc_confusing_routine_name
+ : Warning<"OpenACC 'routine' directive with a name refers to a function "
+ "with the same name as the function on the following line; this "
+ "may be unintended">,
+ InGroup<DiagGroup<"openacc-confusing-routine-name">>;
+def err_acc_decl_for_routine
+ : Error<"expected function or lambda declaration for 'routine' construct">;
// AMDGCN builtins diagnostics
def err_amdgcn_global_load_lds_size_invalid_value : Error<"invalid size value">;
diff --git a/clang/include/clang/Parse/Parser.h b/clang/include/clang/Parse/Parser.h
index fbe2865b1b7c1..c8ceef8f8987d 100644
--- a/clang/include/clang/Parse/Parser.h
+++ b/clang/include/clang/Parse/Parser.h
@@ -3697,7 +3697,10 @@ class Parser : public CodeCompletionHandler {
/// diagnostic. Eventually will be split into a few functions to parse
///
diff erent situations.
public:
- DeclGroupPtrTy ParseOpenACCDirectiveDecl();
+ DeclGroupPtrTy ParseOpenACCDirectiveDecl(AccessSpecifier &AS,
+ ParsedAttributes &Attrs,
+ DeclSpec::TST TagType,
+ Decl *TagDecl);
StmtResult ParseOpenACCDirectiveStmt();
private:
@@ -3831,6 +3834,11 @@ class Parser : public CodeCompletionHandler {
OpenACCGangArgRes ParseOpenACCGangArg(SourceLocation GangLoc);
/// Parses a 'condition' expr, ensuring it results in a
ExprResult ParseOpenACCConditionExpr();
+ DeclGroupPtrTy
+ ParseOpenACCAfterRoutineDecl(AccessSpecifier &AS, ParsedAttributes &Attrs,
+ DeclSpec::TST TagType, Decl *TagDecl,
+ OpenACCDirectiveParseInfo &DirInfo);
+ StmtResult ParseOpenACCAfterRoutineStmt(OpenACCDirectiveParseInfo &DirInfo);
private:
//===--------------------------------------------------------------------===//
diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h
index 358dec4cadb72..3291028c1b621 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -31,6 +31,7 @@
namespace clang {
class IdentifierInfo;
class OpenACCClause;
+class Scope;
class SemaOpenACC : public SemaBase {
public:
@@ -172,6 +173,9 @@ class SemaOpenACC : public SemaBase {
// check them later.
llvm::SmallDenseMap<const clang::FunctionDecl *, SourceLocation>
MagicStaticLocs;
+ OpenACCRoutineDecl *LastRoutineDecl = nullptr;
+
+ void CheckLastRoutineDeclNameConflict(const NamedDecl *ND);
public:
ComputeConstructInfo &getActiveComputeConstructInfo() {
@@ -763,10 +767,37 @@ class SemaOpenACC : public SemaBase {
/// Called after the directive has been completely parsed, including the
/// declaration group or associated statement.
- DeclGroupRef ActOnEndDeclDirective(
- OpenACCDirectiveKind K, SourceLocation StartLoc, SourceLocation DirLoc,
- SourceLocation LParenLoc, Expr *FuncRef, SourceLocation RParenLoc,
- SourceLocation EndLoc, ArrayRef<OpenACCClause *> Clauses);
+ DeclGroupRef
+ ActOnEndDeclDirective(OpenACCDirectiveKind K, SourceLocation StartLoc,
+ SourceLocation DirLoc, SourceLocation LParenLoc,
+ SourceLocation RParenLoc, SourceLocation EndLoc,
+ ArrayRef<OpenACCClause *> Clauses);
+
+ // Helper functions for ActOnEndRoutine*Directive, which does all the checking
+ // given the proper list of declarations.
+ void CheckRoutineDecl(SourceLocation DirLoc,
+ ArrayRef<const OpenACCClause *> Clauses,
+ Decl *NextParsedDecl);
+ OpenACCRoutineDecl *CheckRoutineDecl(SourceLocation StartLoc,
+ SourceLocation DirLoc,
+ SourceLocation LParenLoc, Expr *FuncRef,
+ SourceLocation RParenLoc,
+ ArrayRef<const OpenACCClause *> Clauses,
+ SourceLocation EndLoc);
+ OpenACCRoutineDeclAttr *
+ mergeRoutineDeclAttr(const OpenACCRoutineDeclAttr &Old);
+ DeclGroupRef
+ ActOnEndRoutineDeclDirective(SourceLocation StartLoc, SourceLocation DirLoc,
+ SourceLocation LParenLoc, Expr *ReferencedFunc,
+ SourceLocation RParenLoc,
+ ArrayRef<const OpenACCClause *> Clauses,
+ SourceLocation EndLoc, DeclGroupPtrTy NextDecl);
+ StmtResult
+ ActOnEndRoutineStmtDirective(SourceLocation StartLoc, SourceLocation DirLoc,
+ SourceLocation LParenLoc, Expr *ReferencedFunc,
+ SourceLocation RParenLoc,
+ ArrayRef<const OpenACCClause *> Clauses,
+ SourceLocation EndLoc, Stmt *NextStmt);
/// Called when encountering an 'int-expr' for OpenACC, and manages
/// conversions and diagnostics to 'int'.
@@ -780,8 +811,14 @@ class SemaOpenACC : public SemaBase {
/// Helper function called by ActonVar that is used to check a 'cache' var.
ExprResult ActOnCacheVar(Expr *VarExpr);
/// Function called when a variable declarator is created, which lets us
- /// impelment the 'routine' 'function static variables' restriction.
+ /// implement the 'routine' 'function static variables' restriction.
void ActOnVariableDeclarator(VarDecl *VD);
+ /// Called when a function decl is created, which lets us implement the
+ /// 'routine' 'doesn't match next thing' warning.
+ void ActOnFunctionDeclarator(FunctionDecl *FD);
+ /// Called when a variable is initialized, so we can implement the 'routine
+ /// 'doesn't match the next thing' warning for lambda init.
+ void ActOnVariableInit(VarDecl *VD, QualType InitType);
// Called after 'ActOnVar' specifically for a 'link' clause, which has to do
// some minor additional checks.
diff --git a/clang/include/clang/Serialization/ASTRecordReader.h b/clang/include/clang/Serialization/ASTRecordReader.h
index ae7cd84fbc647..7117b7246739b 100644
--- a/clang/include/clang/Serialization/ASTRecordReader.h
+++ b/clang/include/clang/Serialization/ASTRecordReader.h
@@ -282,6 +282,8 @@ class ASTRecordReader
/// statement reading.
void readOpenACCClauseList(MutableArrayRef<const OpenACCClause *> Clauses);
+ void readOpenACCRoutineDeclAttr(OpenACCRoutineDeclAttr *A);
+
/// Read a source location, advancing Idx.
SourceLocation readSourceLocation(LocSeq *Seq = nullptr) {
return Reader->ReadSourceLocation(*F, Record, Idx, Seq);
diff --git a/clang/include/clang/Serialization/ASTRecordWriter.h b/clang/include/clang/Serialization/ASTRecordWriter.h
index 67720a0aebc1c..84d77e46016b7 100644
--- a/clang/include/clang/Serialization/ASTRecordWriter.h
+++ b/clang/include/clang/Serialization/ASTRecordWriter.h
@@ -307,6 +307,8 @@ class ASTRecordWriter
/// Writes out a list of OpenACC clauses.
void writeOpenACCClauseList(ArrayRef<const OpenACCClause *> Clauses);
+ void AddOpenACCRoutineDeclAttr(const OpenACCRoutineDeclAttr *A);
+
/// Emit a string.
void AddString(StringRef Str) {
return Writer->AddString(Str, *Record);
diff --git a/clang/lib/AST/DeclOpenACC.cpp b/clang/lib/AST/DeclOpenACC.cpp
index f1e6770d83187..760c08d21cccd 100644
--- a/clang/lib/AST/DeclOpenACC.cpp
+++ b/clang/lib/AST/DeclOpenACC.cpp
@@ -12,6 +12,8 @@
#include "clang/AST/DeclOpenACC.h"
#include "clang/AST/ASTContext.h"
+#include "clang/AST/Attr.h"
+#include "clang/AST/OpenACCClause.h"
using namespace clang;
@@ -50,3 +52,12 @@ OpenACCRoutineDecl::CreateDeserialized(ASTContext &Ctx, GlobalDeclID ID,
return new (Ctx, ID, additionalSizeToAlloc<const OpenACCClause *>(NumClauses))
OpenACCRoutineDecl(NumClauses);
}
+
+void OpenACCRoutineDeclAttr::printPrettyPragma(
+ llvm::raw_ostream &OS, const clang::PrintingPolicy &P) const {
+ if (Clauses.size() > 0) {
+ OS << ' ';
+ OpenACCClausePrinter Printer{OS, P};
+ Printer.VisitClauseList(Clauses);
+ }
+}
diff --git a/clang/lib/AST/DeclPrinter.cpp b/clang/lib/AST/DeclPrinter.cpp
index 0b59fa5aec1be..6368531cef3be 100644
--- a/clang/lib/AST/DeclPrinter.cpp
+++ b/clang/lib/AST/DeclPrinter.cpp
@@ -49,6 +49,7 @@ namespace {
QualType T);
void PrintObjCTypeParams(ObjCTypeParamList *Params);
+ void PrintOpenACCRoutineOnLambda(Decl *D);
public:
DeclPrinter(raw_ostream &Out, const PrintingPolicy &Policy,
@@ -292,10 +293,34 @@ bool DeclPrinter::prettyPrintAttributes(const Decl *D,
return hasPrinted;
}
+void DeclPrinter::PrintOpenACCRoutineOnLambda(Decl *D) {
+ CXXRecordDecl *CXXRD = nullptr;
+ if (const auto *VD = dyn_cast<VarDecl>(D)) {
+ if (const auto *Init = VD->getInit())
+ CXXRD = Init->getType().isNull() ? nullptr
+ : Init->getType()->getAsCXXRecordDecl();
+ } else if (const auto *FD = dyn_cast<FieldDecl>(D)) {
+ CXXRD =
+ FD->getType().isNull() ? nullptr : FD->getType()->getAsCXXRecordDecl();
+ }
+
+ if (!CXXRD || !CXXRD->isLambda())
+ return;
+
+ if (const auto *Call = CXXRD->getLambdaCallOperator()) {
+ for (auto *A : Call->specific_attrs<OpenACCRoutineDeclAttr>()) {
+ A->printPretty(Out, Policy);
+ Indent();
+ }
+ }
+}
+
void DeclPrinter::prettyPrintPragmas(Decl *D) {
if (Policy.PolishForDeclaration)
return;
+ PrintOpenACCRoutineOnLambda(D);
+
if (D->hasAttrs()) {
AttrVec &Attrs = D->getAttrs();
for (auto *A : Attrs) {
@@ -894,6 +919,7 @@ void DeclPrinter::VisitFriendDecl(FriendDecl *D) {
}
void DeclPrinter::VisitFieldDecl(FieldDecl *D) {
+ prettyPrintPragmas(D);
// FIXME: add printing of pragma attributes if required.
if (!Policy.SuppressSpecifiers && D->isMutable())
Out << "mutable ";
@@ -1930,19 +1956,17 @@ void DeclPrinter::VisitOpenACCRoutineDecl(OpenACCRoutineDecl *D) {
if (!D->isInvalidDecl()) {
Out << "#pragma acc routine";
- if (D->hasNameSpecified()) {
- Out << "(";
+ Out << "(";
- // The referenced function was named here, but this makes us tolerant of
- // errors.
- if (D->getFunctionReference())
- D->getFunctionReference()->printPretty(Out, nullptr, Policy,
- Indentation, "\n", &Context);
- else
- Out << "<error>";
+ // The referenced function was named here, but this makes us tolerant of
+ // errors.
+ if (D->getFunctionReference())
+ D->getFunctionReference()->printPretty(Out, nullptr, Policy, Indentation,
+ "\n", &Context);
+ else
+ Out << "<error>";
- Out << ")";
- }
+ Out << ")";
if (!D->clauses().empty()) {
Out << ' ';
diff --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp
index 579ee91b7ec99..ab76e6dffa0ff 100644
--- a/clang/lib/AST/OpenACCClause.cpp
+++ b/clang/lib/AST/OpenACCClause.cpp
@@ -628,6 +628,18 @@ OpenACCBindClause *OpenACCBindClause::Create(const ASTContext &C,
return new (Mem) OpenACCBindClause(BeginLoc, LParenLoc, ID, EndLoc);
}
+bool clang::operator==(const OpenACCBindClause &LHS,
+ const OpenACCBindClause &RHS) {
+ if (LHS.isStringArgument() != RHS.isStringArgument())
+ return false;
+
+ if (LHS.isStringArgument())
+ return LHS.getStringArgument()->getString() ==
+ RHS.getStringArgument()->getString();
+ return LHS.getIdentifierArgument()->getName() ==
+ RHS.getIdentifierArgument()->getName();
+}
+
//===----------------------------------------------------------------------===//
// OpenACC clauses printing methods
//===----------------------------------------------------------------------===//
diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp
index 7600b40a7482e..f18cf703bdaa6 100644
--- a/clang/lib/AST/TextNodeDumper.cpp
+++ b/clang/lib/AST/TextNodeDumper.cpp
@@ -3090,10 +3090,7 @@ void TextNodeDumper::VisitOpenACCDeclareDecl(const OpenACCDeclareDecl *D) {
void TextNodeDumper::VisitOpenACCRoutineDecl(const OpenACCRoutineDecl *D) {
OS << " " << D->getDirectiveKind();
- if (D->hasNameSpecified()) {
- OS << " name_specified";
- dumpSourceRange(SourceRange{D->getLParenLoc(), D->getRParenLoc()});
- }
+ dumpSourceRange(SourceRange{D->getLParenLoc(), D->getRParenLoc()});
AddChild([=] { Visit(D->getFunctionReference()); });
@@ -3105,6 +3102,16 @@ void TextNodeDumper::VisitOpenACCRoutineDecl(const OpenACCRoutineDecl *D) {
});
}
+void TextNodeDumper::VisitOpenACCRoutineDeclAttr(
+ const OpenACCRoutineDeclAttr *A) {
+ for (const OpenACCClause *C : A->Clauses)
+ AddChild([=] {
+ Visit(C);
+ for (const Stmt *S : C->children())
+ AddChild([=] { Visit(S); });
+ });
+}
+
void TextNodeDumper::VisitEmbedExpr(const EmbedExpr *S) {
AddChild("begin", [=] { OS << S->getStartingElementPos(); });
AddChild("number of elements", [=] { OS << S->getDataElementCount(); });
diff --git a/clang/lib/Parse/ParseDecl.cpp b/clang/lib/Parse/ParseDecl.cpp
index 9ca3e2b5756ca..3f156407edc99 100644
--- a/clang/lib/Parse/ParseDecl.cpp
+++ b/clang/lib/Parse/ParseDecl.cpp
@@ -5173,7 +5173,9 @@ void Parser::ParseStructUnionBody(SourceLocation RecordLoc,
}
if (Tok.is(tok::annot_pragma_openacc)) {
- ParseOpenACCDirectiveDecl();
+ AccessSpecifier AS = AS_none;
+ ParsedAttributes Attrs(AttrFactory);
+ ParseOpenACCDirectiveDecl(AS, Attrs, TagType, TagDecl);
continue;
}
diff --git a/clang/lib/Parse/ParseDeclCXX.cpp b/clang/lib/Parse/ParseDeclCXX.cpp
index 9384f9ab10af0..51fe0663a8d1a 100644
--- a/clang/lib/Parse/ParseDeclCXX.cpp
+++ b/clang/lib/Parse/ParseDeclCXX.cpp
@@ -3732,7 +3732,7 @@ Parser::DeclGroupPtrTy Parser::ParseCXXClassMemberDeclarationWithPragmas(
return ParseOpenMPDeclarativeDirectiveWithExtDecl(
AS, AccessAttrs, /*Delayed=*/true, TagType, TagDecl);
case tok::annot_pragma_openacc:
- return ParseOpenACCDirectiveDecl();
+ return ParseOpenACCDirectiveDecl(AS, AccessAttrs, TagType, TagDecl);
default:
if (tok::isPragmaAnnotation(Tok.getKind())) {
diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp
index 8d3501082cc27..6c854ddcb505b 100644
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@@ -1542,8 +1542,68 @@ Parser::ParseOpenACCDirective() {
return ParseInfo;
}
+Parser::DeclGroupPtrTy Parser::ParseOpenACCAfterRoutineDecl(
+ AccessSpecifier &AS, ParsedAttributes &Attrs, DeclSpec::TST TagType,
+ Decl *TagDecl, OpenACCDirectiveParseInfo &DirInfo) {
+ assert(DirInfo.DirKind == OpenACCDirectiveKind::Routine);
+
+ DeclGroupPtrTy Ptr;
+ if (DirInfo.LParenLoc.isInvalid()) {
+ if (Tok.isNot(tok::r_brace) && !isEofOrEom()) {
+ if (AS == AS_none) {
+ // This is either an external declaration, or inside of a C struct. If
+ // the latter, we have to diagnose if this is the 'implicit' named
+ // version.
+ if (TagType == DeclSpec::TST_unspecified) {
+ ParsedAttributes EmptyDeclSpecAttrs(AttrFactory);
+ MaybeParseCXX11Attributes(Attrs);
+ ParsingDeclSpec PDS(*this);
+ Ptr = ParseExternalDeclaration(Attrs, EmptyDeclSpecAttrs, &PDS);
+ }
+ // The only way we can have a 'none' access specifier that is in a
+ // not-unspecified tag-type is a C struct. Member functions and
+ // lambdas don't work in C, so we can just count on
+ // ActonRoutineDeclDirective to recognize that Ptr is null and diagnose.
+ } else {
+ Ptr = ParseCXXClassMemberDeclarationWithPragmas(AS, Attrs, TagType,
+ TagDecl);
+ }
+ }
+ }
+
+ return DeclGroupPtrTy::make(
+ getActions().OpenACC().ActOnEndRoutineDeclDirective(
+ DirInfo.StartLoc, DirInfo.DirLoc, DirInfo.LParenLoc,
+ DirInfo.Exprs.empty() ? nullptr : DirInfo.Exprs.front(),
+ DirInfo.RParenLoc, DirInfo.Clauses, DirInfo.EndLoc, Ptr));
+}
+
+StmtResult
+Parser::ParseOpenACCAfterRoutineStmt(OpenACCDirectiveParseInfo &DirInfo) {
+ assert(DirInfo.DirKind == OpenACCDirectiveKind::Routine);
+ // We have to know the next statement for 1 of 2 reasons:
+ // Routine without a name needs an associated DeclStmt.
+ // Routine WITH a name needs to see if it is a DeclStmt to diagnose.
+ StmtResult NextStmt = StmtEmpty();
+
+ // Parse the next statement in the 'implicit' case, not in the 'named' case.
+ // In the 'named' case we will use the creation of the next decl to determine
+ // whether we should warn.
+ if (DirInfo.LParenLoc.isInvalid()) {
+ ParsingOpenACCDirectiveRAII DirScope(*this, /*Value=*/false);
+ NextStmt = ParseStatement();
+ }
+
+ return getActions().OpenACC().ActOnEndRoutineStmtDirective(
+ DirInfo.StartLoc, DirInfo.DirLoc, DirInfo.LParenLoc,
+ DirInfo.Exprs.empty() ? nullptr : DirInfo.Exprs.front(),
+ DirInfo.RParenLoc, DirInfo.Clauses, DirInfo.EndLoc, NextStmt.get());
+}
+
// Parse OpenACC directive on a declaration.
-Parser::DeclGroupPtrTy Parser::ParseOpenACCDirectiveDecl() {
+Parser::DeclGroupPtrTy
+Parser::ParseOpenACCDirectiveDecl(AccessSpecifier &AS, ParsedAttributes &Attrs,
+ DeclSpec::TST TagType, Decl *TagDecl) {
assert(Tok.is(tok::annot_pragma_openacc) && "expected OpenACC Start Token");
ParsingOpenACCDirectiveRAII DirScope(*this);
@@ -1554,9 +1614,11 @@ Parser::DeclGroupPtrTy Parser::ParseOpenACCDirectiveDecl() {
DirInfo.DirKind, DirInfo.StartLoc, DirInfo.Clauses))
return nullptr;
+ if (DirInfo.DirKind == OpenACCDirectiveKind::Routine)
+ return ParseOpenACCAfterRoutineDecl(AS, Attrs, TagType, TagDecl, DirInfo);
+
return DeclGroupPtrTy::make(getActions().OpenACC().ActOnEndDeclDirective(
DirInfo.DirKind, DirInfo.StartLoc, DirInfo.DirLoc, DirInfo.LParenLoc,
- DirInfo.Exprs.empty() ? nullptr : DirInfo.Exprs.front(),
DirInfo.RParenLoc, DirInfo.EndLoc, DirInfo.Clauses));
}
@@ -1571,6 +1633,9 @@ StmtResult Parser::ParseOpenACCDirectiveStmt() {
DirInfo.DirKind, DirInfo.StartLoc, DirInfo.Clauses))
return StmtError();
+ if (DirInfo.DirKind == OpenACCDirectiveKind::Routine)
+ return ParseOpenACCAfterRoutineStmt(DirInfo);
+
StmtResult AssocStmt;
if (doesDirectiveHaveAssociatedStmt(DirInfo.DirKind)) {
SemaOpenACC::AssociatedStmtRAII AssocStmtRAII(
diff --git a/clang/lib/Parse/Parser.cpp b/clang/lib/Parse/Parser.cpp
index 6128bc58d69c0..2eca89179453b 100644
--- a/clang/lib/Parse/Parser.cpp
+++ b/clang/lib/Parse/Parser.cpp
@@ -867,8 +867,11 @@ Parser::ParseExternalDeclaration(ParsedAttributes &Attrs,
AccessSpecifier AS = AS_none;
return ParseOpenMPDeclarativeDirectiveWithExtDecl(AS, Attrs);
}
- case tok::annot_pragma_openacc:
- return ParseOpenACCDirectiveDecl();
+ case tok::annot_pragma_openacc: {
+ AccessSpecifier AS = AS_none;
+ return ParseOpenACCDirectiveDecl(AS, Attrs, DeclSpec::TST_unspecified,
+ /*TagDecl=*/nullptr);
+ }
case tok::annot_pragma_ms_pointers_to_members:
HandlePragmaMSPointersToMembers();
return nullptr;
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 9042655dc1b88..043e82414c052 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -2894,6 +2894,8 @@ static bool mergeDeclAttribute(Sema &S, NamedDecl *D,
else if (isa<SuppressAttr>(Attr))
// Do nothing. Each redeclaration should be suppressed separately.
NewAttr = nullptr;
+ else if (const auto *RD = dyn_cast<OpenACCRoutineDeclAttr>(Attr))
+ NewAttr = S.OpenACC().mergeRoutineDeclAttr(*RD);
else if (Attr->shouldInheritEvenIfAlreadyPresent() || !DeclHasAttr(D, Attr))
NewAttr = cast<InheritableAttr>(Attr->clone(S.Context));
@@ -11028,6 +11030,11 @@ Sema::ActOnFunctionDeclarator(Scope *S, Declarator &D, DeclContext *DC,
D.getFunctionDefinitionKind() == FunctionDefinitionKind::Declaration)
ExternalDeclarations.push_back(NewFD);
+ // Used for a warning on the 'next' declaration when used with a
+ // `routine(name)`.
+ if (getLangOpts().OpenACC)
+ OpenACC().ActOnFunctionDeclarator(NewFD);
+
return NewFD;
}
@@ -14048,6 +14055,9 @@ void Sema::AddInitializerToDecl(Decl *RealDecl, Expr *Init, bool DirectInit) {
VDecl->isFileVarDecl())
DeclsToCheckForDeferredDiags.insert(VDecl);
CheckCompleteVariableDeclaration(VDecl);
+
+ if (LangOpts.OpenACC && !InitType.isNull())
+ OpenACC().ActOnVariableInit(VDecl, InitType);
}
void Sema::ActOnInitializerError(Decl *D) {
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index ec9e9527dca6f..fea334e9e80ff 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -16,6 +16,8 @@
#include "clang/AST/StmtOpenACC.h"
#include "clang/Basic/DiagnosticSema.h"
#include "clang/Basic/OpenACCKinds.h"
+#include "clang/Basic/SourceManager.h"
+#include "clang/Sema/Scope.h"
#include "clang/Sema/Sema.h"
#include "llvm/ADT/StringExtras.h"
#include "llvm/Support/Casting.h"
@@ -1454,9 +1456,11 @@ FunctionDecl *getFunctionFromRoutineName(Expr *RoutineName) {
return nullptr;
RoutineName = RoutineName->IgnoreParenImpCasts();
if (isa<RecoveryExpr>(RoutineName)) {
+ // There is nothing we can do here, this isn't a function we can count on.
return nullptr;
} else if (isa<DependentScopeDeclRefExpr, CXXDependentScopeMemberExpr>(
RoutineName)) {
+ // The lookup is dependent, so we'll have to figure this out later.
return nullptr;
} else if (auto *DRE = dyn_cast<DeclRefExpr>(RoutineName)) {
ValueDecl *VD = DRE->getDecl();
@@ -1466,20 +1470,25 @@ FunctionDecl *getFunctionFromRoutineName(Expr *RoutineName) {
// Allow lambdas.
if (auto *VarD = dyn_cast<VarDecl>(VD)) {
- if (auto *RD = VarD->getType()->getAsCXXRecordDecl()) {
- if (RD->isGenericLambda())
+ QualType VarDTy = VarD->getType();
+ if (!VarDTy.isNull()) {
+ if (auto *RD = VarDTy->getAsCXXRecordDecl()) {
+ if (RD->isGenericLambda())
+ return nullptr;
+ if (RD->isLambda())
+ return RD->getLambdaCallOperator();
+ } else if (VarDTy->isDependentType()) {
+ // We don't really know what this is going to be.
return nullptr;
- if (RD->isLambda())
- return RD->getLambdaCallOperator();
+ }
}
+ return nullptr;
+ } else if (isa<OverloadExpr>(RoutineName)) {
+ return nullptr;
}
- return nullptr;
- } else if (isa<OverloadExpr>(RoutineName)) {
- return nullptr;
}
return nullptr;
}
-
} // namespace
ExprResult SemaOpenACC::ActOnRoutineName(Expr *RoutineName) {
@@ -1502,18 +1511,21 @@ ExprResult SemaOpenACC::ActOnRoutineName(Expr *RoutineName) {
// Allow lambdas.
if (const auto *VarD = dyn_cast<VarDecl>(VD)) {
- if (const auto *RD = VarD->getType()->getAsCXXRecordDecl()) {
- if (RD->isGenericLambda()) {
- Diag(RoutineName->getBeginLoc(), diag::err_acc_routine_overload_set)
- << RoutineName;
- return ExprError();
- }
- if (RD->isLambda())
+ QualType VarDTy = VarD->getType();
+ if (!VarDTy.isNull()) {
+ if (const auto *RD = VarDTy->getAsCXXRecordDecl()) {
+ if (RD->isGenericLambda()) {
+ Diag(RoutineName->getBeginLoc(), diag::err_acc_routine_overload_set)
+ << RoutineName;
+ return ExprError();
+ }
+ if (RD->isLambda())
+ return RoutineName;
+ } else if (VarDTy->isDependentType()) {
+ // If this is a dependent variable, it might be a lambda. So we just
+ // accept this and catch it next time.
return RoutineName;
- } else if (VarD->getType()->isDependentType()) {
- // If this is a dependent variable, it might be a lambda. So we just
- // accept this and catch it next time.
- return RoutineName;
+ }
}
}
@@ -1535,17 +1547,103 @@ ExprResult SemaOpenACC::ActOnRoutineName(Expr *RoutineName) {
return ExprError();
}
void SemaOpenACC::ActOnVariableDeclarator(VarDecl *VD) {
- if (!VD->isStaticLocal())
+ if (!VD->isStaticLocal() || !getLangOpts().OpenACC)
return;
- if (const auto *FD = dyn_cast<FunctionDecl>(getCurContext())) {
- if (const auto *A = FD->getAttr<OpenACCRoutineAnnotAttr>()) {
+ // This cast should be safe, since a static-local can only happen in a
+ // function declaration.
+ auto *ContextDecl = cast<FunctionDecl>(getCurContext());
+
+ // OpenACC 3.3 2.15:
+ // In C and C++, function static variables are not supported in functions to
+ // which a routine directive applies.
+ for (const auto *A : ContextDecl->attrs()) {
+ if (isa<OpenACCRoutineDeclAttr, OpenACCRoutineAnnotAttr>(A)) {
Diag(VD->getBeginLoc(), diag::err_acc_magic_static_in_routine);
Diag(A->getLocation(), diag::note_acc_construct_here)
<< OpenACCDirectiveKind::Routine;
+ return;
}
- MagicStaticLocs.insert({FD, VD->getBeginLoc()});
}
+
+ MagicStaticLocs.insert({ContextDecl->getCanonicalDecl(), VD->getBeginLoc()});
+}
+void SemaOpenACC::CheckLastRoutineDeclNameConflict(const NamedDecl *ND) {
+ // OpenACC 3.3 A.3.4
+ // When a procedure with that name is in scope and it is not the same
+ // procedure as the immediately following procedure declaration or
+ // definition, the resolution of the name can be confusing. Implementations
+ // should then issue a compile-time warning diagnostic even though the
+ // application is conforming.
+
+ // If we haven't created one, also can't diagnose.
+ if (!LastRoutineDecl)
+ return;
+
+ // If the currently created function doesn't have a name, we can't diagnose on
+ // a match.
+ if (!ND->getDeclName().isIdentifier())
+ return;
+
+ // If the two are in
diff erent decl contexts, it doesn't make sense to
+ // diagnose.
+ if (LastRoutineDecl->getDeclContext() != ND->getLexicalDeclContext())
+ return;
+
+ // If we don't have a referenced thing yet, we can't diagnose.
+ FunctionDecl *RoutineTarget =
+ getFunctionFromRoutineName(LastRoutineDecl->getFunctionReference());
+ if (!RoutineTarget)
+ return;
+
+ // If the Routine target doesn't have a name, we can't diagnose.
+ if (!RoutineTarget->getDeclName().isIdentifier())
+ return;
+
+ // Of course don't diagnose if the names don't match.
+ if (ND->getName() != RoutineTarget->getName())
+ return;
+
+ long NDLine = SemaRef.SourceMgr.getSpellingLineNumber(ND->getBeginLoc());
+ long LastLine =
+ SemaRef.SourceMgr.getSpellingLineNumber(LastRoutineDecl->getBeginLoc());
+
+ // Do some line-number math to make sure they are within a line of eachother.
+ // Comments or newlines can be inserted to clarify intent.
+ if (NDLine - LastLine > 1)
+ return;
+
+ // Don't warn if it actually DOES apply to this function via redecls.
+ if (ND->getCanonicalDecl() == RoutineTarget->getCanonicalDecl())
+ return;
+
+ Diag(LastRoutineDecl->getFunctionReference()->getBeginLoc(),
+ diag::warn_acc_confusing_routine_name);
+ Diag(RoutineTarget->getBeginLoc(), diag::note_previous_decl) << ND;
+}
+
+void SemaOpenACC::ActOnVariableInit(VarDecl *VD, QualType InitType) {
+ if (!VD || !getLangOpts().OpenACC || InitType.isNull())
+ return;
+
+ // To avoid double-diagnostic, just diagnose this during instantiation. We'll
+ // get 1 warning per instantiation, but this permits us to be more sensible
+ // for cases where the lookup is confusing.
+ if (VD->getLexicalDeclContext()->isDependentContext())
+ return;
+
+ const auto *RD = InitType->getAsCXXRecordDecl();
+ // If this isn't a lambda, no sense in diagnosing.
+ if (!RD || !RD->isLambda())
+ return;
+
+ CheckLastRoutineDeclNameConflict(VD);
+}
+
+void SemaOpenACC::ActOnFunctionDeclarator(FunctionDecl *FD) {
+ if (!FD || !getLangOpts().OpenACC)
+ return;
+ CheckLastRoutineDeclNameConflict(FD);
}
bool SemaOpenACC::ActOnStartStmtDirective(
@@ -1746,12 +1844,12 @@ StmtResult SemaOpenACC::ActOnEndStmtDirective(
EndLoc);
}
case OpenACCDirectiveKind::Routine:
+ llvm_unreachable("routine shouldn't handled here");
case OpenACCDirectiveKind::Declare: {
// Declare and routine arei declaration directives, but can be used here as
// long as we wrap it in a DeclStmt. So make sure we do that here.
- DeclGroupRef DR = ActOnEndDeclDirective(
- K, StartLoc, DirLoc, LParenLoc,
- (Exprs.empty() ? nullptr : Exprs.front()), RParenLoc, EndLoc, Clauses);
+ DeclGroupRef DR = ActOnEndDeclDirective(K, StartLoc, DirLoc, LParenLoc,
+ RParenLoc, EndLoc, Clauses);
return SemaRef.ActOnDeclStmt(DeclGroupPtrTy::make(DR), StartLoc, EndLoc);
}
@@ -1860,8 +1958,8 @@ bool SemaOpenACC::ActOnStartDeclDirective(
DeclGroupRef SemaOpenACC::ActOnEndDeclDirective(
OpenACCDirectiveKind K, SourceLocation StartLoc, SourceLocation DirLoc,
- SourceLocation LParenLoc, Expr *FuncRef, SourceLocation RParenLoc,
- SourceLocation EndLoc, ArrayRef<OpenACCClause *> Clauses) {
+ SourceLocation LParenLoc, SourceLocation RParenLoc, SourceLocation EndLoc,
+ ArrayRef<OpenACCClause *> Clauses) {
switch (K) {
default:
case OpenACCDirectiveKind::Invalid:
@@ -1882,61 +1980,258 @@ DeclGroupRef SemaOpenACC::ActOnEndDeclDirective(
getCurContext()->addDecl(DeclareDecl);
return DeclGroupRef{DeclareDecl};
}
- case OpenACCDirectiveKind::Routine: {
- // For now, diagnose that we don't support argument-less routine yet.
- if (LParenLoc.isInvalid()) {
- Diag(DirLoc, diag::warn_acc_routine_unimplemented);
- return DeclGroupRef{};
+ case OpenACCDirectiveKind::Routine:
+ llvm_unreachable("routine shouldn't be handled here");
+ }
+ llvm_unreachable("unhandled case in directive handling?");
+}
+
+namespace {
+// Given the decl on the next line, figure out if it is one that is acceptable
+// to `routine`, or looks like the sort of decl we should be diagnosing against.
+FunctionDecl *LegalizeNextParsedDecl(Decl *D) {
+ if (!D)
+ return nullptr;
+
+ // Functions are per-fact acceptable as-is.
+ if (auto *FD = dyn_cast<FunctionDecl>(D))
+ return FD;
+
+ // Function templates are functions, so attach to the templated decl.
+ if (auto *FTD = dyn_cast<FunctionTemplateDecl>(D))
+ return FTD->getTemplatedDecl();
+
+ if (auto *FD = dyn_cast<FieldDecl>(D)) {
+ auto *RD =
+ FD->getType().isNull() ? nullptr : FD->getType()->getAsCXXRecordDecl();
+
+ if (RD && RD->isGenericLambda())
+ return RD->getDependentLambdaCallOperator()->getTemplatedDecl();
+ if (RD && RD->isLambda())
+ return RD->getLambdaCallOperator();
+ }
+ // VarDecl we can look at the init instead of the type of the variable, this
+ // makes us more tolerant of the 'auto' deduced type.
+ if (auto *VD = dyn_cast<VarDecl>(D)) {
+ Expr *Init = VD->getInit();
+ if (!Init || Init->getType().isNull())
+ return nullptr;
+
+ const auto *RD = Init->getType()->getAsCXXRecordDecl();
+ if (RD && RD->isGenericLambda())
+ return RD->getDependentLambdaCallOperator()->getTemplatedDecl();
+ if (RD && RD->isLambda())
+ return RD->getLambdaCallOperator();
+
+ // FIXME: We could try harder in the case where this is a dependent thing
+ // that ends up being a lambda (that is, the init is an unresolved lookup
+ // expr), but we can't attach to the call/lookup expr. If we instead try to
+ // attach to the VarDecl, when we go to instantiate it, attributes are
+ // instantiated before the init, so we can't actually see the type at any
+ // point where it would be relevant/able to be checked. We could perhaps do
+ // some sort of 'after-init' instantiation/checking here, but that doesn't
+ // seem valuable for a situation that other compilers don't handle.
+ }
+ return nullptr;
+}
+
+void CreateRoutineDeclAttr(SemaOpenACC &SemaRef, SourceLocation DirLoc,
+ ArrayRef<const OpenACCClause *> Clauses,
+ ValueDecl *AddTo) {
+ OpenACCRoutineDeclAttr *A =
+ OpenACCRoutineDeclAttr::Create(SemaRef.getASTContext(), DirLoc);
+ A->Clauses.assign(Clauses.begin(), Clauses.end());
+ AddTo->addAttr(A);
+}
+} // namespace
+
+// Variant that adds attributes, because this is the unnamed case.
+void SemaOpenACC::CheckRoutineDecl(SourceLocation DirLoc,
+ ArrayRef<const OpenACCClause *> Clauses,
+ Decl *NextParsedDecl) {
+
+ FunctionDecl *NextParsedFDecl = LegalizeNextParsedDecl(NextParsedDecl);
+
+ if (!NextParsedFDecl) {
+ // If we don't have a valid 'next thing', just diagnose.
+ SemaRef.Diag(DirLoc, diag::err_acc_decl_for_routine);
+ return;
+ }
+
+ // OpenACC 3.3 2.15:
+ // In C and C++, function static variables are not supported in functions to
+ // which a routine directive applies.
+ if (auto Itr = MagicStaticLocs.find(NextParsedFDecl->getCanonicalDecl());
+ Itr != MagicStaticLocs.end()) {
+ Diag(Itr->second, diag::err_acc_magic_static_in_routine);
+ Diag(DirLoc, diag::note_acc_construct_here)
+ << OpenACCDirectiveKind::Routine;
+
+ return;
+ }
+
+ // TODO ERICH: Check bind here.
+ auto BindItr = llvm::find_if(Clauses, llvm::IsaPred<OpenACCBindClause>);
+ for (auto *A : NextParsedFDecl->attrs()) {
+ // OpenACC 3.3 2.15:
+ // If a procedure has a bind clause on both the declaration and definition
+ // than they both must bind to the same name.
+ if (auto *RA = dyn_cast<OpenACCRoutineDeclAttr>(A)) {
+ auto OtherBindItr =
+ llvm::find_if(RA->Clauses, llvm::IsaPred<OpenACCBindClause>);
+ if (OtherBindItr != RA->Clauses.end() &&
+ (*cast<OpenACCBindClause>(*BindItr)) !=
+ (*cast<OpenACCBindClause>(*OtherBindItr))) {
+ Diag((*BindItr)->getBeginLoc(), diag::err_acc_duplicate_unnamed_bind);
+ Diag((*OtherBindItr)->getEndLoc(), diag::note_acc_previous_clause_here);
+ return;
+ }
}
- auto *RoutineDecl = OpenACCRoutineDecl::Create(
- getASTContext(), getCurContext(), StartLoc, DirLoc, LParenLoc, FuncRef,
- RParenLoc, EndLoc, Clauses);
- RoutineDecl->setAccess(AS_public);
- getCurContext()->addDecl(RoutineDecl);
+ // OpenACC 3.3 2.15:
+ // A bind clause may not bind to a routine name that has a visible bind
+ // clause.
+ // We take the combo of these two 2.15 restrictions to mean that the
+ // 'declaration'/'definition' quote is an exception to this. So we're going
+ // to disallow mixing of the two types entirely.
+ if (auto *RA = dyn_cast<OpenACCRoutineAnnotAttr>(A);
+ RA && RA->getRange().getEnd().isValid()) {
+ Diag((*BindItr)->getBeginLoc(), diag::err_acc_duplicate_bind);
+ Diag(RA->getRange().getEnd(), diag::note_acc_previous_clause_here);
+ return;
+ }
+ }
+
+ CreateRoutineDeclAttr(*this, DirLoc, Clauses, NextParsedFDecl);
+}
+
+// Variant that adds a decl, because this is the named case.
+OpenACCRoutineDecl *SemaOpenACC::CheckRoutineDecl(
+ SourceLocation StartLoc, SourceLocation DirLoc, SourceLocation LParenLoc,
+ Expr *FuncRef, SourceLocation RParenLoc,
+ ArrayRef<const OpenACCClause *> Clauses, SourceLocation EndLoc) {
+ assert(LParenLoc.isValid());
+ if (FunctionDecl *FD = getFunctionFromRoutineName(FuncRef)) {
// OpenACC 3.3 2.15:
// In C and C++, function static variables are not supported in functions to
// which a routine directive applies.
- if (auto *FD = getFunctionFromRoutineName(FuncRef)) {
- if (auto Itr = MagicStaticLocs.find(FD); Itr != MagicStaticLocs.end()) {
- Diag(Itr->second, diag::err_acc_magic_static_in_routine);
- Diag(DirLoc, diag::note_acc_construct_here)
- << OpenACCDirectiveKind::Routine;
- }
+ if (auto Itr = MagicStaticLocs.find(FD->getCanonicalDecl());
+ Itr != MagicStaticLocs.end()) {
+ Diag(Itr->second, diag::err_acc_magic_static_in_routine);
+ Diag(DirLoc, diag::note_acc_construct_here)
+ << OpenACCDirectiveKind::Routine;
- // OpenACC 3.3 2.15:
- // A bind clause may not bind to a routine name that has a visible bind
- // clause.
- // TODO OpenACC: There is an exception to this rule that if these are the
- // implicit function style (that is, without a name), they may have
- // duplicates as long as they have the same name.
- auto BindItr = llvm::find_if(Clauses, llvm::IsaPred<OpenACCBindClause>);
- if (auto *A = FD->getAttr<OpenACCRoutineAnnotAttr>()) {
- if (BindItr != Clauses.end()) {
- if (A->BindClause.isInvalid()) {
- // If we have a bind clause, and the function doesn't have one
- // annotated yet, set it.
- A->BindClause = (*BindItr)->getBeginLoc();
- } else {
+ return nullptr;
+ }
+
+ // OpenACC 3.3 2.15:
+ // A bind clause may not bind to a routine name that has a visible bind
+ // clause.
+ auto BindItr = llvm::find_if(Clauses, llvm::IsaPred<OpenACCBindClause>);
+ SourceLocation BindLoc;
+ if (BindItr != Clauses.end()) {
+ BindLoc = (*BindItr)->getBeginLoc();
+ // Since this is adding a 'named' routine, we aren't allowed to combine
+ // with ANY other visible bind clause. Error if we see either.
+
+ for (auto *A : FD->attrs()) {
+ if (auto *RA = dyn_cast<OpenACCRoutineDeclAttr>(A)) {
+ auto OtherBindItr =
+ llvm::find_if(RA->Clauses, llvm::IsaPred<OpenACCBindClause>);
+ if (OtherBindItr != RA->Clauses.end()) {
Diag((*BindItr)->getBeginLoc(), diag::err_acc_duplicate_bind);
- Diag(A->BindClause, diag::note_acc_previous_clause_here);
+ Diag((*OtherBindItr)->getEndLoc(),
+ diag::note_acc_previous_clause_here);
+ return nullptr;
}
}
- } else {
- auto *RAA = OpenACCRoutineAnnotAttr::Create(getASTContext(), DirLoc);
- FD->addAttr(RAA);
- if (BindItr != Clauses.end())
- RAA->BindClause = (*BindItr)->getBeginLoc();
+
+ if (auto *RA = dyn_cast<OpenACCRoutineAnnotAttr>(A);
+ RA && RA->getRange().getEnd().isValid()) {
+ Diag((*BindItr)->getBeginLoc(), diag::err_acc_duplicate_bind);
+ Diag(RA->getRange().getEnd(), diag::note_acc_previous_clause_here);
+ return nullptr;
+ }
}
}
- return DeclGroupRef{RoutineDecl};
+ // Set the end-range to the 'bind' clause here, so we can look it up
+ // later.
+ auto *RAA = OpenACCRoutineAnnotAttr::CreateImplicit(getASTContext(),
+ {DirLoc, BindLoc});
+ FD->addAttr(RAA);
+ // In case we are referencing not the 'latest' version, make sure we add
+ // the attribute to all declarations.
+ while (FD != FD->getMostRecentDecl()) {
+ FD = FD->getMostRecentDecl();
+ FD->addAttr(RAA);
+ }
}
+
+ LastRoutineDecl = OpenACCRoutineDecl::Create(
+ getASTContext(), getCurContext(), StartLoc, DirLoc, LParenLoc, FuncRef,
+ RParenLoc, EndLoc, Clauses);
+ LastRoutineDecl->setAccess(AS_public);
+ getCurContext()->addDecl(LastRoutineDecl);
+
+ return LastRoutineDecl;
+}
+
+DeclGroupRef SemaOpenACC::ActOnEndRoutineDeclDirective(
+ SourceLocation StartLoc, SourceLocation DirLoc, SourceLocation LParenLoc,
+ Expr *ReferencedFunc, SourceLocation RParenLoc,
+ ArrayRef<const OpenACCClause *> Clauses, SourceLocation EndLoc,
+ DeclGroupPtrTy NextDecl) {
+ assert((!ReferencedFunc || !NextDecl) &&
+ "Only one of these should be filled");
+
+ if (LParenLoc.isInvalid()) {
+ Decl *NextLineDecl = nullptr;
+ if (NextDecl && NextDecl.get().isSingleDecl())
+ NextLineDecl = NextDecl.get().getSingleDecl();
+
+ CheckRoutineDecl(DirLoc, Clauses, NextLineDecl);
+
+ return NextDecl.get();
}
- llvm_unreachable("unhandled case in directive handling?");
+
+ return DeclGroupRef{CheckRoutineDecl(
+ StartLoc, DirLoc, LParenLoc, ReferencedFunc, RParenLoc, Clauses, EndLoc)};
}
+StmtResult SemaOpenACC::ActOnEndRoutineStmtDirective(
+ SourceLocation StartLoc, SourceLocation DirLoc, SourceLocation LParenLoc,
+ Expr *ReferencedFunc, SourceLocation RParenLoc,
+ ArrayRef<const OpenACCClause *> Clauses, SourceLocation EndLoc,
+ Stmt *NextStmt) {
+ assert((!ReferencedFunc || !NextStmt) &&
+ "Only one of these should be filled");
+
+ if (LParenLoc.isInvalid()) {
+ Decl *NextLineDecl = nullptr;
+ if (NextStmt)
+ if (DeclStmt *DS = dyn_cast<DeclStmt>(NextStmt); DS && DS->isSingleDecl())
+ NextLineDecl = DS->getSingleDecl();
+
+ CheckRoutineDecl(DirLoc, Clauses, NextLineDecl);
+ return NextStmt;
+ }
+
+ DeclGroupRef DR{CheckRoutineDecl(StartLoc, DirLoc, LParenLoc, ReferencedFunc,
+ RParenLoc, Clauses, EndLoc)};
+ return SemaRef.ActOnDeclStmt(DeclGroupPtrTy::make(DR), StartLoc, EndLoc);
+}
+
+OpenACCRoutineDeclAttr *
+SemaOpenACC::mergeRoutineDeclAttr(const OpenACCRoutineDeclAttr &Old) {
+ OpenACCRoutineDeclAttr *New =
+ OpenACCRoutineDeclAttr::Create(getASTContext(), Old.getLocation());
+ // We should jsut be able to copy these, there isn't really any
+ // merging/inheriting we have to do, so no worry about doing a deep copy.
+ New->Clauses = Old.Clauses;
+ return New;
+}
ExprResult
SemaOpenACC::BuildOpenACCAsteriskSizeExpr(SourceLocation AsteriskLoc) {
return OpenACCAsteriskSizeExpr::Create(getASTContext(), AsteriskLoc);
diff --git a/clang/lib/Sema/SemaOpenACCClause.cpp b/clang/lib/Sema/SemaOpenACCClause.cpp
index 33c2d5d6c8c4c..78e250b759501 100644
--- a/clang/lib/Sema/SemaOpenACCClause.cpp
+++ b/clang/lib/Sema/SemaOpenACCClause.cpp
@@ -1991,6 +1991,9 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitCollapseClause(
OpenACCClause *SemaOpenACCClauseVisitor::VisitBindClause(
SemaOpenACC::OpenACCParsedClause &Clause) {
+ if (checkAlreadyHasClauseOfKind(SemaRef, ExistingClauses, Clause))
+ return nullptr;
+
if (std::holds_alternative<StringLiteral *>(Clause.getBindDetails()))
return OpenACCBindClause::Create(
Ctx, Clause.getBeginLoc(), Clause.getLParenLoc(),
diff --git a/clang/lib/Sema/SemaTemplateInstantiate.cpp b/clang/lib/Sema/SemaTemplateInstantiate.cpp
index 19c27a76b182c..ad49d6f6cce89 100644
--- a/clang/lib/Sema/SemaTemplateInstantiate.cpp
+++ b/clang/lib/Sema/SemaTemplateInstantiate.cpp
@@ -1574,6 +1574,8 @@ namespace {
TransformStmtAlwaysInlineAttr(const Stmt *OrigS, const Stmt *InstS,
const AlwaysInlineAttr *A);
const CodeAlignAttr *TransformCodeAlignAttr(const CodeAlignAttr *CA);
+ const OpenACCRoutineDeclAttr *
+ TransformOpenACCRoutineDeclAttr(const OpenACCRoutineDeclAttr *A);
ExprResult TransformPredefinedExpr(PredefinedExpr *E);
ExprResult TransformDeclRefExpr(DeclRefExpr *E);
ExprResult TransformCXXDefaultArgExpr(CXXDefaultArgExpr *E);
@@ -2253,6 +2255,12 @@ TemplateInstantiator::TransformCodeAlignAttr(const CodeAlignAttr *CA) {
Expr *TransformedExpr = getDerived().TransformExpr(CA->getAlignment()).get();
return getSema().BuildCodeAlignAttr(*CA, TransformedExpr);
}
+const OpenACCRoutineDeclAttr *
+TemplateInstantiator::TransformOpenACCRoutineDeclAttr(
+ const OpenACCRoutineDeclAttr *A) {
+ llvm_unreachable("RoutineDecl should only be a declaration attribute, as it "
+ "applies to a Function Decl (and a few places for VarDecl)");
+}
ExprResult TemplateInstantiator::transformNonTypeTemplateParmRef(
Decl *AssociatedDecl, const NonTypeTemplateParmDecl *parm,
diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
index 539c2fdb83797..8aaaea0bcdd66 100644
--- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
+++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
@@ -769,6 +769,11 @@ attrToRetainOwnershipKind(const Attr *A) {
}
}
+// Implementation is down with the rest of the OpenACC Decl instantiations.
+static void instantiateDependentOpenACCRoutineDeclAttr(
+ Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs,
+ const OpenACCRoutineDeclAttr *OldAttr, const Decl *Old, Decl *New);
+
void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs,
const Decl *Tmpl, Decl *New,
LateInstantiatedAttrVec *LateAttrs,
@@ -868,6 +873,12 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs,
continue;
}
+ if (const auto *RoutineAttr = dyn_cast<OpenACCRoutineDeclAttr>(TmplAttr)) {
+ instantiateDependentOpenACCRoutineDeclAttr(*this, TemplateArgs,
+ RoutineAttr, Tmpl, New);
+ continue;
+ }
+
// Existing DLL attribute on the instantiation takes precedence.
if (TmplAttr->getKind() == attr::DLLExport ||
TmplAttr->getKind() == attr::DLLImport) {
@@ -1319,6 +1330,25 @@ llvm::SmallVector<OpenACCClause *> InstantiateOpenACCClauseList(
} // namespace
+static void instantiateDependentOpenACCRoutineDeclAttr(
+ Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs,
+ const OpenACCRoutineDeclAttr *OldAttr, const Decl *OldDecl, Decl *NewDecl) {
+ OpenACCRoutineDeclAttr *A =
+ OpenACCRoutineDeclAttr::Create(S.getASTContext(), OldAttr->getLocation());
+
+ if (!OldAttr->Clauses.empty()) {
+ llvm::SmallVector<OpenACCClause *> TransformedClauses =
+ InstantiateOpenACCClauseList(
+ S, TemplateArgs, OpenACCDirectiveKind::Routine, OldAttr->Clauses);
+ A->Clauses.assign(TransformedClauses.begin(), TransformedClauses.end());
+ }
+
+ // We don't end up having to do any magic-static or bind checking here, since
+ // the first phase should have caught this, since we always apply to the
+ // functiondecl.
+ NewDecl->addAttr(A);
+}
+
Decl *TemplateDeclInstantiator::VisitOpenACCDeclareDecl(OpenACCDeclareDecl *D) {
SemaRef.OpenACC().ActOnConstruct(D->getDirectiveKind(), D->getBeginLoc());
llvm::SmallVector<OpenACCClause *> TransformedClauses =
@@ -1330,8 +1360,8 @@ Decl *TemplateDeclInstantiator::VisitOpenACCDeclareDecl(OpenACCDeclareDecl *D) {
return nullptr;
DeclGroupRef Res = SemaRef.OpenACC().ActOnEndDeclDirective(
- D->getDirectiveKind(), D->getBeginLoc(), D->getDirectiveLoc(), {},
- nullptr, {}, D->getEndLoc(), TransformedClauses);
+ D->getDirectiveKind(), D->getBeginLoc(), D->getDirectiveLoc(), {}, {},
+ D->getEndLoc(), TransformedClauses);
if (Res.isNull())
return nullptr;
@@ -1358,10 +1388,9 @@ Decl *TemplateDeclInstantiator::VisitOpenACCRoutineDecl(OpenACCRoutineDecl *D) {
D->getDirectiveKind(), D->getBeginLoc(), TransformedClauses))
return nullptr;
- DeclGroupRef Res = SemaRef.OpenACC().ActOnEndDeclDirective(
- D->getDirectiveKind(), D->getBeginLoc(), D->getDirectiveLoc(),
- D->getLParenLoc(), FuncRef.get(), D->getRParenLoc(), D->getEndLoc(),
- TransformedClauses);
+ DeclGroupRef Res = SemaRef.OpenACC().ActOnEndRoutineDeclDirective(
+ D->getBeginLoc(), D->getDirectiveLoc(), D->getLParenLoc(), FuncRef.get(),
+ D->getRParenLoc(), TransformedClauses, D->getEndLoc(), nullptr);
if (Res.isNull())
return nullptr;
@@ -1671,6 +1700,9 @@ Decl *TemplateDeclInstantiator::VisitVarDecl(VarDecl *D,
if (Var->getTLSKind())
SemaRef.CheckThreadLocalForLargeAlignment(Var);
+ if (SemaRef.getLangOpts().OpenACC)
+ SemaRef.OpenACC().ActOnVariableDeclarator(Var);
+
return Var;
}
diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
index c66ea5a2c0346..a07d34a49550d 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -12871,6 +12871,12 @@ void ASTRecordReader::readOpenACCClauseList(
Clauses[I] = readOpenACCClause();
}
+void ASTRecordReader::readOpenACCRoutineDeclAttr(OpenACCRoutineDeclAttr *A) {
+ unsigned NumVars = readInt();
+ A->Clauses.resize(NumVars);
+ readOpenACCClauseList(A->Clauses);
+}
+
static unsigned getStableHashForModuleName(StringRef PrimaryModuleName) {
// TODO: Maybe it is better to check PrimaryModuleName is a valid
// module name?
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index 436d5213c70dc..95064f3a36f2c 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -8871,3 +8871,10 @@ void ASTRecordWriter::writeOpenACCClauseList(
for (const OpenACCClause *Clause : Clauses)
writeOpenACCClause(Clause);
}
+void ASTRecordWriter::AddOpenACCRoutineDeclAttr(
+ const OpenACCRoutineDeclAttr *A) {
+ // We have to write the size so that the reader can do a resize. Unlike the
+ // Decl version of this, we can't count on trailing storage to get this right.
+ writeUInt32(A->Clauses.size());
+ writeOpenACCClauseList(A->Clauses);
+}
diff --git a/clang/test/AST/ast-print-openacc-routine-construct.cpp b/clang/test/AST/ast-print-openacc-routine-construct.cpp
index 1d21e2ddee9a1..b3de3402a952c 100644
--- a/clang/test/AST/ast-print-openacc-routine-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-routine-construct.cpp
@@ -1,85 +1,143 @@
// RUN: %clang_cc1 -fopenacc -ast-print %s -o - | FileCheck %s
auto Lambda = [](){};
-// CHECK: #pragma acc routine(Lambda) worker bind(identifier)
+// CHECK: auto Lambda = []() {
#pragma acc routine(Lambda) worker bind(identifier)
+// CHECK: #pragma acc routine(Lambda) worker bind(identifier)
+
int function();
-// CHECK: #pragma acc routine(function) vector nohost bind("string")
#pragma acc routine (function) vector nohost bind("string")
-
-// CHECK: #pragma acc routine(function) device_type(Something) seq
+// CHECK: #pragma acc routine(function) vector nohost bind("string")
#pragma acc routine(function) device_type(Something) seq
-// CHECK: #pragma acc routine(function) dtype(Something) seq
+// CHECK-NEXT: #pragma acc routine(function) device_type(Something) seq
#pragma acc routine(function) dtype(Something) seq
+// CHECK-NEXT: #pragma acc routine(function) dtype(Something) seq
+
+#pragma acc routine nohost vector
+int function2();
+// CHECK: #pragma acc routine nohost vector
+// CHECK-NEXT: int function2()
+
+#pragma acc routine worker nohost bind("asdf")
+auto Lambda2 = [](){};
+// CHECK: #pragma acc routine worker nohost bind("asdf")
+// CHECK-NEXT: auto Lambda2 = []() {
+#pragma acc routine worker nohost bind("asdf")
+auto Lambda3 = [](auto){};
+// CHECK: #pragma acc routine worker nohost bind("asdf")
+// CHECK-NEXT: auto Lambda3 = [](auto) {
namespace NS {
int NSFunc();
auto Lambda = [](){};
}
-// CHECK: #pragma acc routine(NS::NSFunc) seq
+
#pragma acc routine(NS::NSFunc) seq
-// CHECK: #pragma acc routine(NS::Lambda) nohost gang
+// CHECK: #pragma acc routine(NS::NSFunc) seq
#pragma acc routine(NS::Lambda) nohost gang
constexpr int getInt() { return 1; }
struct S {
+ // CHECK: struct S {
+ // despite being targetted by 'named' versions, we shouldn't print the
+ // attribute here.
+ // CHECK-NEXT: void MemFunc();
void MemFunc();
+#pragma acc routine gang(dim: 1)
+ void MemFunc2();
+// CHECK-NEXT: #pragma acc routine gang(dim: 1)
+// CHECK-NEXT: void MemFunc2();
static void StaticMemFunc();
+// CHECK-NEXT: static void StaticMemFunc();
+#pragma acc routine gang(dim: getInt())
+ static void StaticMemFunc2();
+// CHECK-NEXT: #pragma acc routine gang(dim: getInt())
+// CHECK-NEXT: static void StaticMemFunc2();
+
constexpr static auto Lambda = [](){};
-// CHECK: #pragma acc routine(S::MemFunc) gang(dim: 1)
+// CHECK-NEXT: static constexpr auto Lambda = []() {
+#pragma acc routine worker
+ constexpr static auto Lambda2 = [](){ return 1; };
+// CHECK: #pragma acc routine worker
+// CHECK-NEXT: static constexpr auto Lambda2 = []() {
+
#pragma acc routine(S::MemFunc) gang(dim:1)
-// CHECK: #pragma acc routine(S::StaticMemFunc) gang(dim: getInt())
+// CHECK: #pragma acc routine(S::MemFunc) gang(dim: 1)
#pragma acc routine(S::StaticMemFunc) gang(dim:getInt())
-// CHECK: #pragma acc routine(S::Lambda) worker
+// CHECK-NEXT: #pragma acc routine(S::StaticMemFunc) gang(dim: getInt())
#pragma acc routine(S::Lambda) worker
+// CHECK-NEXT: #pragma acc routine(S::Lambda) worker
-// CHECK: #pragma acc routine(MemFunc) gang(dim: 1)
#pragma acc routine(MemFunc) gang(dim:1)
-// CHECK: #pragma acc routine(StaticMemFunc) gang(dim: getInt())
+// CHECK-NEXT: #pragma acc routine(MemFunc) gang(dim: 1)
#pragma acc routine(StaticMemFunc) gang(dim:getInt())
-// CHECK: #pragma acc routine(Lambda) nohost worker
+// CHECK-NEXT: #pragma acc routine(StaticMemFunc) gang(dim: getInt())
#pragma acc routine(Lambda) nohost worker
+// CHECK-NEXT: #pragma acc routine(Lambda) nohost worker
};
-// CHECK: #pragma acc routine(S::MemFunc) gang(dim: 1)
#pragma acc routine(S::MemFunc) gang(dim:1)
-// CHECK: #pragma acc routine(S::StaticMemFunc) worker
+// CHECK: #pragma acc routine(S::MemFunc) gang(dim: 1)
#pragma acc routine(S::StaticMemFunc) worker
-// CHECK: #pragma acc routine(S::Lambda) vector
+// CHECK-NEXT: #pragma acc routine(S::StaticMemFunc) worker
#pragma acc routine(S::Lambda) vector
+// CHECK-NEXT: #pragma acc routine(S::Lambda) vector
template<typename T>
struct DepS {
void MemFunc();
+// CHECK: void MemFunc();
+
static void StaticMemFunc();
+// CHECK-NEXT: static void StaticMemFunc();
+
+#pragma acc routine gang(dim: T{1})
+ static T StaticMemFunc2();
+// CHECK-NEXT: #pragma acc routine gang(dim: T{1})
+// CHECK-NEXT: static T StaticMemFunc2();
+
constexpr static auto Lambda = [](){ return 1;};
+// CHECK-NEXT: static constexpr auto Lambda = []() {
+
+#pragma acc routine gang(dim: T{1})
+ constexpr static auto Lambda2 = [](){return 1;};
+// CHECK: #pragma acc routine gang(dim: T{1})
+// CHECK-NEXT: static constexpr auto Lambda2 = []() {
+#pragma acc routine gang(dim: T{1})
+ constexpr static auto Lambda3 = [](auto){return 1;};
+// CHECK: #pragma acc routine gang(dim: T{1})
+// CHECK-NEXT: static constexpr auto Lambda3 = [](auto) {
+#pragma acc routine gang(dim: Lambda())
+ T MemFunc2();
+// CHECK: #pragma acc routine gang(dim: Lambda())
+// CHECK-NEXT: T MemFunc2();
-// CHECK: #pragma acc routine(Lambda) gang(dim: Lambda())
#pragma acc routine(Lambda) gang(dim:Lambda())
-// CHECK: #pragma acc routine(MemFunc) worker
+// CHECK-NEXT: #pragma acc routine(Lambda) gang(dim: Lambda())
#pragma acc routine(MemFunc) worker
-// CHECK: #pragma acc routine(StaticMemFunc) seq
+// CHECK-NEXT: #pragma acc routine(MemFunc) worker
#pragma acc routine(StaticMemFunc) seq
+// CHECK-NEXT: #pragma acc routine(StaticMemFunc) seq
-// 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
+// CHECK-NEXT: #pragma acc routine(DepS<T>::Lambda) gang(dim: 1)
#pragma acc routine(DepS::MemFunc) gang
-// CHECK: #pragma acc routine(DepS<T>::StaticMemFunc) worker
+// CHECK-NEXT: #pragma acc routine(DepS<T>::MemFunc) gang
#pragma acc routine(DepS::StaticMemFunc) worker
+// CHECK-NEXT: #pragma acc routine(DepS<T>::StaticMemFunc) worker
-// CHECK: #pragma acc routine(DepS<T>::Lambda) vector
#pragma acc routine(DepS<T>::Lambda) vector
-// CHECK: #pragma acc routine(DepS<T>::MemFunc) seq nohost
+// CHECK-NEXT: #pragma acc routine(DepS<T>::Lambda) vector
#pragma acc routine(DepS<T>::MemFunc) seq nohost
-// CHECK: #pragma acc routine(DepS<T>::StaticMemFunc) nohost worker
+// CHECK-NEXT: #pragma acc routine(DepS<T>::MemFunc) seq nohost
#pragma acc routine(DepS<T>::StaticMemFunc) nohost worker
+// CHECK-NEXT: #pragma acc routine(DepS<T>::StaticMemFunc) nohost worker
-// CHECK: #pragma acc routine(MemFunc) worker dtype(*)
#pragma acc routine (MemFunc) worker dtype(*)
-// CHECK: #pragma acc routine(MemFunc) device_type(Lambda) vector
+// CHECK-NEXT: #pragma acc routine(MemFunc) worker dtype(*)
#pragma acc routine (MemFunc) device_type(Lambda) vector
+// CHECK-NEXT: #pragma acc routine(MemFunc) device_type(Lambda) vector
};
// CHECK: #pragma acc routine(DepS<int>::Lambda) gang bind("string")
@@ -98,4 +156,17 @@ void TemplFunc() {
#pragma acc routine(T::StaticMemFunc) worker nohost bind(identifier)
// CHECK: #pragma acc routine(T::Lambda) nohost seq bind("string")
#pragma acc routine(T::Lambda) nohost seq bind("string")
+
+ auto Lambda1 = [](){};
+#pragma acc routine(Lambda1) seq
+// CHECK: #pragma acc routine(Lambda1) seq
+#pragma acc routine seq
+ auto Lambda2 = [](){};
+// CHECK: #pragma acc routine seq
+// CHECK-NEXT: auto Lambda2 = []() {
+#pragma acc routine seq
+ auto Lambda3 = [](auto){};
+// CHECK: #pragma acc routine seq
+// CHECK-NEXT: auto Lambda3 = [](auto) {
+ Lambda3(T{});
}
diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c
index 5178138216773..100dc07a9fd60 100644
--- a/clang/test/ParserOpenACC/parse-clauses.c
+++ b/clang/test/ParserOpenACC/parse-clauses.c
@@ -1261,7 +1261,6 @@ void Gang() {
}
- // 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}}
@@ -1278,15 +1277,14 @@ void bar();
// Bind Clause Parsing.
- // expected-error at +2{{expected '('}}
- // expected-warning at +1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}}
+// expected-error at +1{{expected '('}}
#pragma acc routine seq bind
void BCP1();
// expected-error at +1{{expected identifier or string literal}}
#pragma acc routine(BCP1) seq bind()
- // expected-warning at +1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}}
+// expected-error at +1{{expected function or lambda declaration for 'routine' construct}}
#pragma acc routine seq bind("ReductionClauseParsing")
#pragma acc routine(BCP1) seq bind(unknown_thing)
diff --git a/clang/test/ParserOpenACC/parse-constructs.c b/clang/test/ParserOpenACC/parse-constructs.c
index 969c31b5bd8c1..71e47abd3a825 100644
--- a/clang/test/ParserOpenACC/parse-constructs.c
+++ b/clang/test/ParserOpenACC/parse-constructs.c
@@ -150,7 +150,6 @@ void func() {
for(;;){}
}
-// expected-warning at +1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}}
#pragma acc routine seq
void routine_func();
// expected-error at +2{{invalid OpenACC clause 'clause'}}
diff --git a/clang/test/ParserOpenACC/parse-constructs.cpp b/clang/test/ParserOpenACC/parse-constructs.cpp
index 2c083f558ab12..814f6a1fd09f2 100644
--- a/clang/test/ParserOpenACC/parse-constructs.cpp
+++ b/clang/test/ParserOpenACC/parse-constructs.cpp
@@ -51,3 +51,10 @@ namespace NS {
// expected-note@#PrivateMemFunc{{implicitly declared private here}}
#pragma acc routine (NS::C::private_mem_func) seq
#pragma acc routine (NS::C::public_mem_func) seq
+
+void foo() {
+ auto x = [](){};
+#pragma acc routine seq
+ auto y = [](){};
+#pragma acc routine (x) seq
+}
diff --git a/clang/test/SemaOpenACC/routine-construct-ast.cpp b/clang/test/SemaOpenACC/routine-construct-ast.cpp
index 7715ec62e7d4e..068d99123522e 100644
--- a/clang/test/SemaOpenACC/routine-construct-ast.cpp
+++ b/clang/test/SemaOpenACC/routine-construct-ast.cpp
@@ -6,46 +6,98 @@
#ifndef PCH_HELPER
#define PCH_HELPER
+
auto Lambda = [](){};
#pragma acc routine(Lambda) worker nohost bind("string")
-// CHECK: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' '(lambda at
// CHECK-NEXT: worker clause
// CHECK-NEXT: nohost clause
// CHECK-NEXT: bind clause
// CHECK-NEXT: StringLiteral{{.*}} "string"
+
int function();
#pragma acc routine (function) nohost vector bind(identifier)
-// CHECK: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'function' 'int ()'
// CHECK-NEXT: nohost clause
// CHECK-NEXT: vector clause
// CHECK-NEXT: bind clause identifier 'identifier'
#pragma acc routine(function) device_type(Something) seq
-// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'function' 'int ()'
// CHECK-NEXT: device_type(Something)
// CHECK-NEXT: seq clause
#pragma acc routine(function) nohost dtype(Something) vector
-// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'function' 'int ()'
// CHECK-NEXT: nohost clause
// CHECK-NEXT: dtype(Something)
// CHECK-NEXT: vector clause
+#pragma acc routine nohost vector
+int function2();
+// CHECK: FunctionDecl{{.*}} function2
+// CHECK-NEXT: OpenACCRoutineDeclAttr
+// CHECK-NEXT: nohost clause
+// CHECK-NEXT: vector clause
+
+#pragma acc routine worker nohost bind("asdf")
+auto Lambda2 = [](){};
+// CHECK: VarDecl{{.*}} Lambda2 '(lambda at
+// CHECK: CXXMethodDecl{{.*}} operator()
+// CHECK-NEXT: CompoundStmt
+// CHECK-NEXT: OpenACCRoutineDeclAttr
+// CHECK-NEXT: worker clause
+// CHECK-NEXT: nohost clause
+// CHECK-NEXT: bind clause
+// CHECK-NEXT: StringLiteral{{.*}}"asdf"
+
+#pragma acc routine worker nohost bind("asdf")
+auto Lambda3 = [](auto a){ return a;};
+// CHECK: VarDecl{{.*}} Lambda3 '(lambda at
+// CHECK-NEXT: LambdaExpr
+// CHECK-NEXT: CXXRecordDecl
+// CHECK: FunctionTemplateDecl{{.*}} operator()
+// CHECK-NEXT: TemplateTypeParmDecl
+// CHECK-NEXT: CXXMethodDecl{{.*}} operator() 'auto (auto) const'
+// CHECK-NEXT: ParmVarDecl{{.*}}'auto'
+// CHECK-NEXT: CompoundStmt
+// CHECK-NEXT: ReturnStmt
+// CHECK-NEXT: DeclRefExpr
+// CHECK: OpenACCRoutineDeclAttr
+// CHECK-NEXT: worker clause
+// CHECK-NEXT: nohost clause
+// CHECK-NEXT: bind clause
+// CHECK-NEXT: StringLiteral{{.*}}"asdf"
+// CHECK-NEXT: CXXMethodDecl{{.*}}operator() 'int (int) const'{{.*}} implicit_instantiation
+// CHECK-NEXT: TemplateArgument type 'int'
+// CHECK-NEXT: BuiltinType
+// CHECK-NEXT: ParmVarDecl{{.*}}'int'
+// CHECK-NEXT: CompoundStmt
+// CHECK-NEXT: ReturnStmt
+// CHECK-NEXT: ImplicitCastExpr
+// CHECK-NEXT: DeclRefExpr
+// CHECK-NEXT: OpenACCRoutineDeclAttr
+// CHECK-NEXT: worker clause
+// CHECK-NEXT: nohost clause
+// CHECK-NEXT: bind clause
+// CHECK-NEXT: StringLiteral{{.*}}"asdf"
+
+int InstL3 = Lambda3(1);
+
namespace NS {
- // CHECK-NEXT: NamespaceDecl
int NSFunc();
auto Lambda = [](){};
}
#pragma acc routine(NS::NSFunc) seq
-// CHECK: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'NSFunc' 'int ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'NS'
// CHECK-NEXT: seq clause
#pragma acc routine(NS::Lambda) gang
-// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'NS::(lambda at
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'NS'
// CHECK-NEXT: gang clause
@@ -54,70 +106,98 @@ constexpr int getInt() { return 1; }
struct S {
void MemFunc();
+#pragma acc routine gang(dim: 1)
+ void MemFunc2();
+ // CHECK: CXXMethodDecl{{.*}}MemFunc2
+ // CHECK-NEXT: OpenACCRoutineDeclAttr{{.*}}
+ // CHECK-NEXT: gang clause
+ // CHECK-NEXT: ConstantExpr{{.*}}'int'
+ // CHECK-NEXT: value: Int 1
static void StaticMemFunc();
+#pragma acc routine gang(dim: getInt())
+ static void StaticMemFunc2();
+ // CHECK: CXXMethodDecl{{.*}}StaticMemFunc2
+ // CHECK-NEXT: OpenACCRoutineDeclAttr{{.*}}
+ // CHECK-NEXT: gang clause
+ // CHECK-NEXT: ConstantExpr{{.*}}'int'
+ // CHECK-NEXT: value: Int 1
constexpr static auto Lambda = [](){ return 1; };
+#pragma acc routine worker
+ constexpr static auto Lambda2 = [](){ return 1; };
+ // CHECK: VarDecl{{.*}}Lambda2 'const S::(lambda at
+ // CHECK-NEXT: value: Struct
+ // CHECK-NEXT: ImplicitCastExpr
+ // CHECK-NEXT: LambdaExpr
+ // CHECK-NEXT: CXXRecordDecl
+ // CHECK: CXXMethodDecl{{.*}} operator() 'int () const'
+ // CHECK-NEXT: CompoundStmt
+ // CHECK-NEXT: ReturnStmt
+ // CHECK-NEXT: IntegerLiteral
+ // CHECK: OpenACCRoutineDeclAttr{{.*}}
+ // CHECK-NEXT: worker clause
+
#pragma acc routine(S::MemFunc) gang(dim: 1)
-// CHECK: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'void ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'S'
// 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: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'void ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'S'
// 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: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const S::(lambda at
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'S'
// CHECK-NEXT: worker clause
#pragma acc routine(MemFunc) gang(dim: 1)
-// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'void ()'
// 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: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'void ()'
// 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: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const S::(lambda at
// CHECK-NEXT: worker clause
#pragma acc routine(Lambda) worker device_type(Lambda)
-// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const S::(lambda at
// CHECK-NEXT: worker clause
// CHECK-NEXT: device_type(Lambda)
#pragma acc routine(Lambda) dtype(Lambda) vector
-// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const S::(lambda at
// CHECK-NEXT: dtype(Lambda)
// CHECK-NEXT: vector clause
};
#pragma acc routine(S::MemFunc) gang(dim: 1)
-// CHECK: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'void ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'S'
// 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: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'void ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'S'
// CHECK-NEXT: worker clause
#pragma acc routine(S::Lambda) vector
-// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const S::(lambda at
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'S'
// CHECK-NEXT: vector
@@ -126,57 +206,96 @@ template<typename T>
struct DepS {
T MemFunc();
static T StaticMemFunc();
+#pragma acc routine gang(dim: T{1})
+ static T StaticMemFunc2();
+ // CHECK: CXXMethodDecl{{.*}}StaticMemFunc2
+ // CHECK-NEXT: OpenACCRoutineDeclAttr{{.*}}
+ // CHECK-NEXT: gang clause
+ // CHECK-NEXT: CXXUnresolvedConstructExpr{{.*}}'T' list
constexpr static auto Lambda = [](){return 1;};
+#pragma acc routine gang(dim: T{1})
+ constexpr static auto Lambda2 = [](){return 1;};
+ // CHECK: VarDecl{{.*}}Lambda2 'const auto'
+ // CHECK-NEXT: LambdaExpr
+ // CHECK-NEXT: CXXRecordDecl
+ // CHECK: CXXMethodDecl{{.*}} operator()
+ // CHECK-NEXT: CompoundStmt
+ // CHECK-NEXT: ReturnStmt
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1
+ // CHECK-NEXT: OpenACCRoutineDeclAttr{{.*}}
+ // CHECK-NEXT: gang clause
+ // CHECK-NEXT: CXXUnresolvedConstructExpr{{.*}}'T' list
+#pragma acc routine gang(dim: T{1})
+ constexpr static auto Lambda3 = [](auto){return 1;};
+ // CHECK: VarDecl{{.*}}Lambda3 'const auto'
+ // CHECK-NEXT: LambdaExpr
+ // CHECK-NEXT: CXXRecordDecl
+ // CHECK: FunctionTemplateDecl{{.*}} operator()
+ // CHECK-NEXT: TemplateTypeParmDecl
+ // CHECK-NEXT: CXXMethodDecl{{.*}} operator() 'auto (auto) const'
+ // CHECK-NEXT: ParmVarDecl{{.*}}'auto'
+ // CHECK-NEXT: CompoundStmt
+ // CHECK-NEXT: ReturnStmt
+ // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1
+ // CHECK-NEXT: OpenACCRoutineDeclAttr{{.*}}
+ // CHECK-NEXT: gang clause
+ // CHECK-NEXT: CXXUnresolvedConstructExpr{{.*}}'T' list
+#pragma acc routine gang(dim: Lambda())
+ T MemFunc2();
+ // CHECK: CXXMethodDecl{{.*}}MemFunc2
+ // CHECK-NEXT: OpenACCRoutineDeclAttr{{.*}}
+ // CHECK-NEXT: gang clause
+ // CHECK-NEXT: CallExpr{{.*}}'<dependent type>'
#pragma acc routine(Lambda) gang(dim: Lambda())
-// CHECK: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const auto'
// CHECK-NEXT: gang clause
// CHECK-NEXT: CallExpr{{.*}}'<dependent type>'
#pragma acc routine(MemFunc) worker
-// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'T ()'
// CHECK-NEXT: worker clause
#pragma acc routine(StaticMemFunc) seq
-// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'T ()'
// CHECK-NEXT: seq clause
#pragma acc routine(DepS::Lambda) gang(dim:1)
-// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const auto'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<T>'
// 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: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'T ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<T>'
// CHECK-NEXT: gang clause
#pragma acc routine(DepS::StaticMemFunc) worker
-// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'T ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<T>'
// CHECK-NEXT: worker clause
#pragma acc routine(DepS<T>::Lambda) vector
-// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const auto'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<T>'
// CHECK-NEXT: vector clause
#pragma acc routine(DepS<T>::MemFunc) seq
-// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'T ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<T>'
// CHECK-NEXT: seq clause
#pragma acc routine(DepS<T>::StaticMemFunc) worker
-// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'T ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<T>'
// CHECK-NEXT: worker clause
#pragma acc routine(DepS<T>::StaticMemFunc) worker device_type(T)
-// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'T ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<T>'
// CHECK-NEXT: worker clause
@@ -186,53 +305,96 @@ struct DepS {
// CHECK: ClassTemplateSpecializationDecl{{.*}}struct DepS
// CHECK: CXXRecordDecl{{.*}} struct DepS
-// CHECK: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK: CXXMethodDecl{{.*}}StaticMemFunc2
+// CHECK-NEXT: OpenACCRoutineDeclAttr{{.*}}
+// CHECK-NEXT: gang clause
+// CHECK-NEXT: ConstantExpr{{.*}}'int'
+// CHECK-NEXT: value: Int 1
+
+// CHECK: VarDecl{{.*}}Lambda2 'const DepS<int>::(lambda at
+// CHECK-NEXT: value: Struct
+// CHECK-NEXT: ImplicitCastExpr
+// CHECK-NEXT: LambdaExpr
+// CHECK-NEXT: CXXRecordDecl
+// CHECK: CXXMethodDecl{{.*}} operator()
+// CHECK-NEXT: CompoundStmt
+// CHECK-NEXT: ReturnStmt
+// CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1
+// CHECK-NEXT: OpenACCRoutineDeclAttr{{.*}}
+// CHECK-NEXT: gang clause
+// CHECK-NEXT: ConstantExpr{{.*}}'int'
+// CHECK-NEXT: value: Int 1
+
+// CHECK: VarDecl{{.*}}Lambda3 'const DepS<int>::(lambda at
+// CHECK-NEXT: value: Struct
+// CHECK-NEXT: ImplicitCastExpr
+// CHECK-NEXT: LambdaExpr
+// CHECK-NEXT: CXXRecordDecl
+// CHECK: FunctionTemplateDecl{{.*}} operator()
+// CHECK-NEXT: TemplateTypeParmDecl
+// CHECK-NEXT: CXXMethodDecl{{.*}} operator() 'auto (auto) const'
+// CHECK-NEXT: ParmVarDecl{{.*}}'auto'
+// CHECK-NEXT: CompoundStmt
+// CHECK-NEXT: ReturnStmt
+// CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1
+// CHECK-NEXT: OpenACCRoutineDeclAttr{{.*}}
+// CHECK-NEXT: gang clause
+// CHECK-NEXT: ConstantExpr{{.*}}'int'
+// CHECK-NEXT: value: Int 1
+
+// CHECK: CXXMethodDecl{{.*}}MemFunc2
+// CHECK-NEXT: OpenACCRoutineDeclAttr{{.*}}
+// CHECK-NEXT: gang clause
+// CHECK-NEXT: ConstantExpr{{.*}}'int'
+// CHECK-NEXT: value: Int 1
+
+// CHECK: OpenACCRoutineDecl{{.*}} routine
// 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: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'int ()'
// CHECK-NEXT: worker clause
-// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'int ()'
// CHECK-NEXT: seq clause
-// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
// 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: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'int ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>'
// CHECK-NEXT: gang clause
-// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'int ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>'
// CHECK-NEXT: worker clause
-// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const DepS<int>::
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>'
// CHECK-NEXT: vector clause
-// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'int ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>'
// CHECK-NEXT: seq clause
-// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'int ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>'
// CHECK-NEXT: worker clause
-// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'int ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>'
// CHECK-NEXT: worker clause
@@ -240,19 +402,19 @@ struct DepS {
};
#pragma acc routine(DepS<int>::Lambda) gang(dim:1)
-// CHECK: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK: OpenACCRoutineDecl{{.*}} routine
// 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
#pragma acc routine(DepS<int>::MemFunc) worker
-// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'int ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>'
// CHECK-NEXT: worker clause
#pragma acc routine(DepS<int>::StaticMemFunc) vector
-// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'int ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>'
// CHECK-NEXT: vector clause
@@ -261,14 +423,14 @@ template<typename T>
void TemplFunc() {
#pragma acc routine(T::MemFunc) gang(dim:T::Lambda())
// CHECK: DeclStmt
-// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DependentScopeDeclRefExpr{{.*}}'<dependent type>'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'T'
// CHECK-NEXT: gang clause
// CHECK-NEXT: CallExpr{{.*}}'<dependent type>'
#pragma acc routine(T::StaticMemFunc) nohost worker bind("string")
// CHECK-NEXT: DeclStmt
-// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DependentScopeDeclRefExpr{{.*}}'<dependent type>'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'T'
// CHECK-NEXT: nohost clause
@@ -277,17 +439,48 @@ void TemplFunc() {
// CHECK-NEXT: StringLiteral{{.*}} "string"
#pragma acc routine(T::Lambda) seq nohost bind(identifier)
// CHECK-NEXT: DeclStmt
-// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine
// CHECK-NEXT: DependentScopeDeclRefExpr{{.*}}'<dependent type>'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'T'
// CHECK-NEXT: seq clause
// CHECK-NEXT: nohost clause
// CHECK-NEXT: bind clause identifier 'identifier'
+ auto Lambda1 = [](){};
+#pragma acc routine(Lambda1) seq
+// CHECK: OpenACCRoutineDecl{{.*}} routine
+// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda1' 'auto'
+// CHECK-NEXT: seq clause
+
+#pragma acc routine seq
+ auto Lambda2 = [](){};
+ // CHECK: VarDecl{{.*}} Lambda2 'auto'
+ // CHECK-NEXT: LambdaExpr
+ // CHECK-NEXT: CXXRecordDecl
+ // CHECK: CXXMethodDecl{{.*}} operator()
+ // CHECK-NEXT: CompoundStmt
+ // CHECK-NEXT: OpenACCRoutineDeclAttr{{.*}}
+ // CHECK-NEXT: seq clause
+
+#pragma acc routine seq
+ auto Lambda3 = [](auto){};
+ // CHECK: VarDecl{{.*}} Lambda3 'auto'
+ // CHECK-NEXT: LambdaExpr
+ // CHECK-NEXT: CXXRecordDecl
+ // CHECK: FunctionTemplateDecl{{.*}} operator()
+ // CHECK-NEXT: TemplateTypeParmDecl
+ // CHECK-NEXT: CXXMethodDecl{{.*}} operator() 'auto (auto) const'
+ // CHECK-NEXT: ParmVarDecl{{.*}}'auto'
+ // CHECK-NEXT: CompoundStmt
+ // CHECK-NEXT: OpenACCRoutineDeclAttr
+ // CHECK-NEXT: seq clause
+
+ Lambda3(T{});
+
// Instantiation:
// CHECK: FunctionDecl{{.*}} TemplFunc 'void ()' implicit_instantiation
// CHECK: DeclStmt
-// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}}
// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'void ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'S'
// CHECK-NEXT: gang clause
@@ -295,7 +488,7 @@ void TemplFunc() {
// CHECK-NEXT: value: Int 1
// CHECK-NEXT: DeclStmt
-// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}}
// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'void ()'
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'S'
// CHECK-NEXT: nohost clause
@@ -304,17 +497,48 @@ void TemplFunc() {
// CHECK-NEXT: StringLiteral{{.*}} "string"
// CHECK-NEXT: DeclStmt
-// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified
+// CHECK-NEXT: OpenACCRoutineDecl{{.*}}
// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const S::(lambda at
// CHECK-NEXT: NestedNameSpecifier{{.*}} 'S'
// CHECK-NEXT: seq clause
// CHECK-NEXT: nohost clause
// CHECK-NEXT: bind clause identifier 'identifier'
+
+// CHECK: OpenACCRoutineDecl{{.*}} routine
+// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda1' '(lambda at
+// CHECK-NEXT: seq clause
+
+// CHECK: VarDecl{{.*}} Lambda2 '(lambda at
+// CHECK-NEXT: LambdaExpr
+// CHECK-NEXT: CXXRecordDecl
+// CHECK: CXXMethodDecl{{.*}} operator()
+// CHECK-NEXT: CompoundStmt
+// CHECK-NEXT: OpenACCRoutineDeclAttr{{.*}}
+// CHECK-NEXT: seq clause
+
+// CHECK: VarDecl{{.*}} Lambda3 '(lambda at
+// CHECK-NEXT: LambdaExpr
+// CHECK-NEXT: CXXRecordDecl
+// CHECK: FunctionTemplateDecl{{.*}} operator()
+// CHECK-NEXT: TemplateTypeParmDecl
+// CHECK-NEXT: CXXMethodDecl{{.*}} operator() 'auto (auto) const'
+// CHECK-NEXT: ParmVarDecl{{.*}}'auto'
+// CHECK-NEXT: CompoundStmt
+// CHECK-NEXT: OpenACCRoutineDeclAttr
+// CHECK-NEXT: seq clause
+// CHECK: CXXMethodDecl{{.*}} operator() 'void (S) const'
+// CHECK-NEXT: TemplateArgument
+// CHECK-NEXT: RecordType
+// CHECK-NEXT: CXXRecord
+// CHECK-NEXT: ParmVarDecl
+// CHECK-NEXT: CompoundStmt
+// CHECK-NEXT: OpenACCRoutineDeclAttr
+// CHECK-NEXT: seq clause
}
void usage() {
DepS<int> s;
TemplFunc<S>();
}
-#endif
+#endif
diff --git a/clang/test/SemaOpenACC/routine-construct-clauses.cpp b/clang/test/SemaOpenACC/routine-construct-clauses.cpp
index da8ed382c4d63..78554fa70292a 100644
--- a/clang/test/SemaOpenACC/routine-construct-clauses.cpp
+++ b/clang/test/SemaOpenACC/routine-construct-clauses.cpp
@@ -7,6 +7,14 @@ void Func2();
#pragma acc routine(Func) vector nohost
#pragma acc routine(Func) nohost seq
#pragma acc routine(Func) gang
+// expected-error at +2{{OpenACC 'bind' clause cannot appear more than once on a 'routine' directive}}
+// expected-note at +1{{previous clause is here}}
+#pragma acc routine(Func) gang bind(a) bind(a)
+
+// expected-error at +2{{OpenACC 'bind' clause cannot appear more than once on a 'routine' directive}}
+// expected-note at +1{{previous clause is here}}
+#pragma acc routine gang bind(a) bind(a)
+void DupeImplName();
// 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}}
@@ -117,6 +125,14 @@ struct DependentT {
#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 MemFunc();
+// expected-error at +1{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 4}}
+#pragma acc routine(MemFunc) gang(dim:T::Four())
+
+// expected-error at +1{{argument to 'gang' clause dimension must be 1, 2, or 3: evaluated to 4}}
+#pragma acc routine gang(dim:T::Four())
+ void MemFunc2();
};
void Inst() {
@@ -132,8 +148,9 @@ void Inst() {
#pragma acc routine(Func) device_type(*) worker
#pragma acc routine(Func) device_type(*) vector
#pragma acc routine(Func) dtype(*) seq
-#pragma acc routine(Func) seq device_type(*) bind("asdf")
-#pragma acc routine(Func2) seq device_type(*) bind(WhateverElse)
+#pragma acc routine(Func2) seq device_type(*) bind("asdf")
+void Func6();
+#pragma acc routine(Func6) seq device_type(*) bind(WhateverElse)
#pragma acc routine(Func) seq dtype(*) device_type(*)
// expected-error at +2{{OpenACC clause 'nohost' may not follow a 'dtype' clause in a 'routine' construct}}
// expected-note at +1{{previous clause is here}}
@@ -152,8 +169,8 @@ void Func3();
#pragma acc routine(Func3) seq bind("asdf")
void Func4();
-#pragma acc routine(Func4) seq
// OK: Doesn't have a bind
+#pragma acc routine(Func4) seq
#pragma acc routine(Func4) seq bind("asdf")
// expected-error at +2{{multiple 'routine' directives with 'bind' clauses are not permitted to refer to the same function}}
// expected-note at -2{{previous clause is here}}
@@ -165,3 +182,48 @@ void Func5();
#pragma acc routine(Func5) seq bind("asdf")
// OK: Doesn't have a bind
#pragma acc routine(Func5) seq
+
+// OK, same.
+#pragma acc routine bind("asdf") seq
+void DupeBinds1();
+#pragma acc routine bind("asdf") seq
+void DupeBinds1();
+
+#pragma acc routine bind(asdf) seq
+void DupeBinds2();
+#pragma acc routine bind(asdf) seq
+void DupeBinds2();
+
+#pragma acc routine bind("asdf") seq
+void DupeBinds3();
+// expected-error at +2{{OpenACC 'bind' clause on a declaration must bind to the same name as previous bind clauses}}
+// expected-note at -3{{previous clause is here}}
+#pragma acc routine bind(asdf) seq
+void DupeBinds3();
+
+void DupeBinds4();
+#pragma acc routine(DupeBinds4) bind(asdf) seq
+// expected-error at +2{{multiple 'routine' directives with 'bind' clauses are not permitted to refer to the same function}}
+// expected-note at -2{{previous clause is here}}
+#pragma acc routine bind(asdf) seq
+void DupeBinds4();
+
+void DupeBinds4b();
+// expected-error at +1{{expected function or lambda declaration for 'routine' construct}}
+#pragma acc routine bind(asdf) seq
+#pragma acc routine(DupeBinds4b) bind(asdf) seq
+void DupeBinds4b();
+
+#pragma acc routine bind(asdf) seq
+void DupeBinds5();
+// expected-error at +2{{OpenACC 'bind' clause on a declaration must bind to the same name as previous bind clauses}}
+// expected-note at -3{{previous clause is here}}
+#pragma acc routine bind(asdfDiff) seq
+void DupeBinds5();
+
+#pragma acc routine bind("asdf") seq
+void DupeBinds6();
+// expected-error at +2{{OpenACC 'bind' clause on a declaration must bind to the same name as previous bind clauses}}
+// expected-note at -3{{previous clause is here}}
+#pragma acc routine bind("asdfDiff") seq
+void DupeBinds6();
diff --git a/clang/test/SemaOpenACC/routine-construct.cpp b/clang/test/SemaOpenACC/routine-construct.cpp
index 3c7a31a2c0774..170cafe50ebdb 100644
--- a/clang/test/SemaOpenACC/routine-construct.cpp
+++ b/clang/test/SemaOpenACC/routine-construct.cpp
@@ -11,6 +11,14 @@ void SameFunc();
#pragma acc routine(SameFunc) seq
void SameFunc();
+namespace NS {
+void DifferentFunc();
+};
+// expected-warning at +2{{OpenACC 'routine' directive with a name refers to a function with the same name as the function on the following line; this may be unintended}}
+// expected-note at -3{{'DifferentFunc' declared here}}
+#pragma acc routine(NS::DifferentFunc) seq
+void DifferentFunc();
+
void NoMagicStatic() {
static int F = 1;
}
@@ -26,6 +34,21 @@ void NoMagicStatic2() {
static int F = 1;
}
+#pragma acc routine seq
+void NoMagicStatic3() {
+ // expected-error at +2{{function static variables are not permitted in functions to which an OpenACC 'routine' directive applies}}
+ // expected-note at -3{{'routine' construct is here}}
+ static int F = 1;
+}
+
+#pragma acc routine seq
+void NoMagicStatic4();
+void NoMagicStatic4() {
+ // expected-error at +2{{function static variables are not permitted in functions to which an OpenACC 'routine' directive applies}}
+ // expected-note at -4{{'routine' construct is here}}
+ static int F = 1;
+}
+
void HasMagicStaticLambda() {
auto MSLambda = []() {
static int I = 5;
@@ -33,16 +56,48 @@ 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) seq
+
+// 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 seq
+ auto MSLambda2 = []() {
+ static int I = 5;
+ };
+
+// Properly handle error recovery.
+// expected-error at +1{{expected function or lambda declaration for 'routine' construct}}
+#pragma acc routine seq
+ auto MSLambda2 = [](auto) {
+ // expected-error at -1{{redefinition of 'MSLambda2'}}
+ // expected-note at -9{{previous definition is here}}
+ static int I = 5;
+ };
+// 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 seq
+ auto MSLambda3 = [](auto) {
+ static int I = 5;
+ };
}
auto Lambda = [](){};
#pragma acc routine(Lambda) seq
+
+#pragma acc routine seq
+auto Lambda2 = [](){};
auto GenLambda = [](auto){};
// expected-error at +1{{OpenACC routine name 'GenLambda' names a set of overloads}}
#pragma acc routine(GenLambda) seq
+
+#pragma acc routine seq
+auto GenLambda2 = [](auto){};
+
// Variable?
+// expected-error at +1{{expected function or lambda declaration for 'routine' construct}}
+#pragma acc routine seq
int Variable;
// Plain function
+#pragma acc routine seq
int function();
#pragma acc routine (function) seq
@@ -50,6 +105,8 @@ int function();
#pragma acc routine (Variable) seq
// Var template?
+// expected-error at +1{{expected function or lambda declaration for 'routine' construct}}
+#pragma acc routine seq
template<typename T>
T VarTempl = 0;
// expected-error at +2{{use of variable template 'VarTempl' requires template arguments}}
@@ -72,6 +129,11 @@ int ambig_func(int);
// expected-error at +1{{OpenACC routine name 'ambig_func' names a set of overloads}}
#pragma acc routine (ambig_func) seq
+#pragma acc routine seq
+int ambig_func2();
+#pragma acc routine seq
+int ambig_func2(int);
+
// Ambiguous in NS
namespace NS {
int ambig_func();
@@ -83,6 +145,9 @@ int ambig_func(int);
// function template
template<typename T, typename U>
void templ_func();
+#pragma acc routine seq
+template<typename T, typename U>
+void templ_func2();
// expected-error at +1{{OpenACC routine name 'templ_func' names a set of overloads}}
#pragma acc routine(templ_func) seq
@@ -93,11 +158,21 @@ void templ_func();
struct S {
void MemFunc();
+#pragma acc routine seq
+ void MemFunc2();
static void StaticMemFunc();
+#pragma acc routine seq
+ static void StaticMemFunc2();
template<typename U>
void TemplMemFunc();
+#pragma acc routine seq
+ template<typename U>
+ void TemplMemFunc2();
template<typename U>
static void TemplStaticMemFunc();
+#pragma acc routine seq
+ template<typename U>
+ static void TemplStaticMemFunc2();
void MemFuncAmbig();
void MemFuncAmbig(int);
@@ -109,6 +184,10 @@ struct S {
int Field;
constexpr static auto Lambda = [](){};
+#pragma acc routine seq
+ constexpr static auto Lambda2 = [](){};
+#pragma acc routine seq
+ constexpr static auto Lambda3 = [](auto){};
#pragma acc routine(S::MemFunc) seq
#pragma acc routine(S::StaticMemFunc) seq
@@ -151,6 +230,21 @@ struct S {
// expected-error at +1{{OpenACC routine name 'S::Field' does not name a function}}
#pragma acc routine(S::Field) seq
+constexpr auto getLambda() {
+ return [](){};
+}
+template<typename T>
+constexpr auto getTemplLambda() {
+ return [](T){};
+}
+constexpr auto getDepLambda() {
+ return [](auto){};
+}
+template<typename T>
+constexpr auto getTemplDepLambda() {
+ return [](auto){};
+}
+
template<typename T>
struct DepS { // #DEPS
void MemFunc();
@@ -174,6 +268,37 @@ struct DepS { // #DEPS
// expected-note@#DEPSInst{{in instantiation of template class}}
static auto LambdaBroken = [](){};
+#pragma acc routine seq
+ constexpr static auto LambdaKinda = getLambda();
+ // FIXME: We can't really handle this/things like this, see comment in
+ // SemaOpenACC.cpp's LegalizeNextParsedDecl.
+ // expected-error at +1{{expected function or lambda declaration for 'routine' construct}}
+#pragma acc routine seq
+ constexpr static auto LambdaKinda2 = getTemplLambda<T>();
+#pragma acc routine seq
+ constexpr static auto DepLambdaKinda = getDepLambda();
+ // expected-error at +1{{expected function or lambda declaration for 'routine' construct}}
+#pragma acc routine seq
+ constexpr static auto DepLambdaKinda2 = getTemplDepLambda<T>();
+
+// expected-error at +1{{expected function or lambda declaration for 'routine' construct}}
+#pragma acc routine seq
+ constexpr static auto Bad = T{};
+
+#pragma acc routine seq
+ constexpr static auto LambdaHasMagicStatic = []() {
+ // expected-error at +2{{function static variables are not permitted in functions to which an OpenACC 'routine' directive applies}}
+ // expected-note at -3{{'routine' construct is here}}
+ static int F = 1;
+ };
+
+ void HasMagicStatic() {
+ static int F = 1; // #HasMagicStaticFunc
+ }
+ void HasMagicStatic2() {
+ static int F = 1; // #HasMagicStaticFunc2
+ }
+
#pragma acc routine(DepS::MemFunc) seq
#pragma acc routine(DepS::StaticMemFunc) seq
#pragma acc routine(DepS::Lambda) seq
@@ -215,10 +340,38 @@ struct DepS { // #DEPS
#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) seq
+
+// FIXME: We could do better about suppressing this double diagnostic, but we
+// don't want to invalidate the vardecl for openacc, so we don't have a good
+// way to do this in the AST.
+// expected-error@#HasMagicStaticFunc 2{{function static variables are not permitted in functions to which an OpenACC 'routine' directive applies}}
+// expected-note at +2 2{{'routine' construct is here}}
+// expected-note at +1{{in instantiation of member function}}}
+#pragma acc routine(DepS<T>::HasMagicStatic) seq
};
+template<typename T>
+void DepF() {
+#pragma acc routine seq
+ auto LambdaKinda = getLambda();
+// expected-error at +1{{expected function or lambda declaration for 'routine' construct}}
+#pragma acc routine seq
+ auto LambdaKinda2 = getTemplLambda<T>();
+#pragma acc routine seq
+ auto DepLambdaKinda = getDepLambda();
+// expected-error at +1{{expected function or lambda declaration for 'routine' construct}}
+#pragma acc routine seq
+ auto DepLambdaKinda2 = getTemplDepLambda<T>();
+
+// expected-error at +1{{expected function or lambda declaration for 'routine' construct}}
+#pragma acc routine seq
+ constexpr static auto Bad = T{};
+}
+
void Inst() {
DepS<int> S; // #DEPSInst
+ S.HasMagicStatic2();
+ DepF<int>(); // #DEPFInst
}
@@ -300,6 +453,11 @@ void Inst() {
// expected-error at +1{{OpenACC routine name 'DepS<int>::Field' does not name a function}}
#pragma acc routine(DepS<int>::Field) seq
+// expected-error@#HasMagicStaticFunc2{{function static variables are not permitted in functions to which an OpenACC 'routine' directive applies}}
+// expected-note at +2{{'routine' construct is here}}
+// expected-note at +1{{in instantiation of member function}}}
+#pragma acc routine(DepS<int>::HasMagicStatic2) seq
+
template<typename T>
void TemplFunc() {
#pragma acc routine(T::MemFunc) seq
@@ -388,6 +546,13 @@ struct DepRefersToT {
#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) seq
+
+// expected-error at +1{{expected function or lambda declaration for 'routine' construct}}
+#pragma acc routine seq
+ auto L = getTemplLambda<U>();
+// expected-error at +1{{expected function or lambda declaration for 'routine' construct}}
+#pragma acc routine seq
+ auto L2 = getTemplDepLambda<U>();
}
};
@@ -398,3 +563,60 @@ void inst() {
s.MemFunc(); // expected-note{{in instantiation of}}
s.TemplMemFunc<S>(); // expected-note{{in instantiation of}}
}
+
+// A.3.4 tests:
+
+void DiffFuncs(); // #GLOBALDIFFFUNCS
+namespace NS {
+// expected-warning at +2{{OpenACC 'routine' directive with a name refers to a function with the same name as the function on the following line; this may be unintended}}
+// expected-note@#GLOBALDIFFFUNCS{{'DiffFuncs' declared here}}
+#pragma acc routine(DiffFuncs) seq
+void DiffFuncs();
+}
+
+void has_
diff _func() {
+// expected-warning at +2{{OpenACC 'routine' directive with a name refers to a function with the same name as the function on the following line; this may be unintended}}
+// expected-note@#GLOBALDIFFFUNCS{{'DiffFuncs' declared here}}
+#pragma acc routine(DiffFuncs) seq
+auto DiffFuncs = [](){};
+}
+
+template<typename T>
+void has_
diff _func_templ() {
+// expected-warning at +3{{OpenACC 'routine' directive with a name refers to a function with the same name as the function on the following line; this may be unintended}}
+// expected-note@#GLOBALDIFFFUNCS{{'DiffFuncs' declared here}}
+// expected-note@#HDFT_INST{{in instantiation of function template specialization}}
+#pragma acc routine(DiffFuncs) seq
+auto DiffFuncs = [](){};
+}
+
+void inst_
diff () {
+ has_
diff _func_templ<int>();// #HDFT_INST
+}
+
+struct SDiff {
+// expected-warning at +2{{OpenACC 'routine' directive with a name refers to a function with the same name as the function on the following line; this may be unintended}}
+// expected-note@#GLOBALDIFFFUNCS{{'DiffFuncs' declared here}}
+#pragma acc routine(DiffFuncs) seq
+ void DiffFuncs();
+};
+template<typename T>
+struct TemplSDiff {
+// expected-warning at +2{{OpenACC 'routine' directive with a name refers to a function with the same name as the function on the following line; this may be unintended}}
+// expected-note@#GLOBALDIFFFUNCS{{'DiffFuncs' declared here}}
+#pragma acc routine(DiffFuncs) seq
+ void DiffFuncs();
+};
+
+struct SOperator {
+#pragma acc routine(DiffFuncs) seq
+ bool operator==(const S&);
+};
+
+namespace NS2 {
+ // Shouldn't diagnose.
+#pragma acc routine(DiffFuncs) seq
+#pragma acc routine seq
+ void DiffFuncs();
+};
+
diff --git a/clang/test/SemaOpenACC/unimplemented-construct.c b/clang/test/SemaOpenACC/unimplemented-construct.c
deleted file mode 100644
index 4e29189de13b6..0000000000000
--- a/clang/test/SemaOpenACC/unimplemented-construct.c
+++ /dev/null
@@ -1,38 +0,0 @@
-// RUN: %clang_cc1 %s -verify -fopenacc
-
-// expected-warning at +1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}}
-#pragma acc routine seq
-
-struct S {
-// expected-warning at +1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}}
-#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 seq
- int foo;
-
-// expected-warning at +1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}}
-#pragma acc routine seq
- {
-// expected-warning at +1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}}
-#pragma acc routine seq
- {
-// expected-warning at +1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}}
-#pragma acc routine seq
- }
- }
-
-// expected-warning at +1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}}
-#pragma acc routine seq
- while(0){}
-
-// expected-warning at +1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}}
-#pragma acc routine seq
- for(;;){}
-
-// expected-warning at +1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}}
-#pragma acc routine seq
-};
diff --git a/clang/utils/TableGen/ClangAttrEmitter.cpp b/clang/utils/TableGen/ClangAttrEmitter.cpp
index 3d90bfbc75b32..06f6b073240ba 100644
--- a/clang/utils/TableGen/ClangAttrEmitter.cpp
+++ b/clang/utils/TableGen/ClangAttrEmitter.cpp
@@ -3533,6 +3533,11 @@ void EmitClangAttrPCHRead(const RecordKeeper &Records, raw_ostream &OS) {
DelayedArgs->writePCHReadArgs(OS);
OS << ");\n";
}
+
+ if (Attr->getValueAsBit("HasCustomSerialization"))
+ OS << " read" << R.getName() << "Attr(cast<" << R.getName()
+ << "Attr>(New));\n";
+
OS << " break;\n";
OS << " }\n";
}
@@ -3563,6 +3568,10 @@ void EmitClangAttrPCHWrite(const RecordKeeper &Records, raw_ostream &OS) {
for (const auto *Arg : Args)
createArgument(*Arg, R.getName())->writePCHWrite(OS);
+
+ if (Attr->getValueAsBit("HasCustomSerialization"))
+ OS << " Record.Add" << R.getName() << "Attr(SA);\n";
+
OS << " break;\n";
OS << " }\n";
}
More information about the cfe-commits
mailing list