diff options
| author | erichkeane <ekeane@nvidia.com> | 2025-01-22 06:47:14 -0800 |
|---|---|---|
| committer | erichkeane <ekeane@nvidia.com> | 2025-03-06 06:42:17 -0800 |
| commit | df1e102e2a6b0e0f1ecf28c58a4a51dbcbe74360 (patch) | |
| tree | dd051896f103fc9c9567dbd680ffff507b8fac54 | |
| parent | 4703f8b6610a13b549c1b8aefe90d0f8975fea1e (diff) | |
| download | llvm-df1e102e2a6b0e0f1ecf28c58a4a51dbcbe74360.zip llvm-df1e102e2a6b0e0f1ecf28c58a4a51dbcbe74360.tar.gz llvm-df1e102e2a6b0e0f1ecf28c58a4a51dbcbe74360.tar.bz2 | |
[OpenACC] implement AST/Sema for 'routine' construct with argument
The 'routine' construct has two forms, one which takes the name of a
function that it applies to, and another where it implicitly figures it
out based on the next declaration. This patch implements the former with
the required restrictions on the name and the function-static-variables
as specified.
What has not been implemented is any clauses for this, any of the A.3.4
warnings, or the other form.
37 files changed, 1191 insertions, 187 deletions
diff --git a/clang/include/clang/AST/DeclOpenACC.h b/clang/include/clang/AST/DeclOpenACC.h index fdeb131..9e99061 100644 --- a/clang/include/clang/AST/DeclOpenACC.h +++ b/clang/include/clang/AST/DeclOpenACC.h @@ -101,6 +101,61 @@ public: static bool classof(const Decl *D) { return classofKind(D->getKind()); } static bool classofKind(Kind K) { return K == OpenACCDeclare; } }; + +class OpenACCRoutineDecl final + : public OpenACCConstructDecl, + private llvm::TrailingObjects<OpenACCRoutineDecl, const OpenACCClause *> { + friend TrailingObjects; + friend class ASTDeclReader; + friend class ASTDeclWriter; + + Expr *FuncRef = nullptr; + SourceRange ParensLoc; + + OpenACCRoutineDecl(unsigned NumClauses) + : OpenACCConstructDecl(OpenACCRoutine) { + std::uninitialized_value_construct( + getTrailingObjects<const OpenACCClause *>(), + getTrailingObjects<const OpenACCClause *>() + NumClauses); + setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(), + NumClauses)); + } + + OpenACCRoutineDecl(DeclContext *DC, SourceLocation StartLoc, + SourceLocation DirLoc, SourceLocation LParenLoc, + Expr *FuncRef, SourceLocation RParenLoc, + SourceLocation EndLoc, + ArrayRef<const OpenACCClause *> Clauses) + : OpenACCConstructDecl(OpenACCRoutine, DC, OpenACCDirectiveKind::Routine, + StartLoc, DirLoc, EndLoc), + FuncRef(FuncRef), ParensLoc(LParenLoc, RParenLoc) { + // Initialize the trailing storage. + std::uninitialized_copy(Clauses.begin(), Clauses.end(), + getTrailingObjects<const OpenACCClause *>()); + setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(), + Clauses.size())); + } + +public: + static OpenACCRoutineDecl * + Create(ASTContext &Ctx, DeclContext *DC, SourceLocation StartLoc, + SourceLocation DirLoc, SourceLocation LParenLoc, Expr *FuncRef, + SourceLocation RParenLoc, SourceLocation EndLoc, + ArrayRef<const OpenACCClause *> Clauses); + static OpenACCRoutineDecl * + CreateDeserialized(ASTContext &Ctx, GlobalDeclID ID, unsigned NumClauses); + static bool classof(const Decl *D) { return classofKind(D->getKind()); } + 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 #endif diff --git a/clang/include/clang/AST/JSONNodeDumper.h b/clang/include/clang/AST/JSONNodeDumper.h index 97b07e1..570662b 100644 --- a/clang/include/clang/AST/JSONNodeDumper.h +++ b/clang/include/clang/AST/JSONNodeDumper.h @@ -282,6 +282,7 @@ public: void VisitBlockDecl(const BlockDecl *D); void VisitOpenACCDeclareDecl(const OpenACCDeclareDecl *D); + void VisitOpenACCRoutineDecl(const OpenACCRoutineDecl *D); void VisitDeclRefExpr(const DeclRefExpr *DRE); void VisitSYCLUniqueStableNameExpr(const SYCLUniqueStableNameExpr *E); diff --git a/clang/include/clang/AST/RecursiveASTVisitor.h b/clang/include/clang/AST/RecursiveASTVisitor.h index 89757dd..5ca3e43 100644 --- a/clang/include/clang/AST/RecursiveASTVisitor.h +++ b/clang/include/clang/AST/RecursiveASTVisitor.h @@ -1825,6 +1825,11 @@ DEF_TRAVERSE_DECL(OMPAllocateDecl, { DEF_TRAVERSE_DECL(OpenACCDeclareDecl, { TRY_TO(VisitOpenACCClauseList(D->clauses())); }) +DEF_TRAVERSE_DECL(OpenACCRoutineDecl, { + TRY_TO(TraverseStmt(D->getFunctionReference())); + TRY_TO(VisitOpenACCClauseList(D->clauses())); +}) + // A helper method for TemplateDecl's children. template <typename Derived> bool RecursiveASTVisitor<Derived>::TraverseTemplateParameterListHelper( diff --git a/clang/include/clang/AST/TextNodeDumper.h b/clang/include/clang/AST/TextNodeDumper.h index 0925aca..ad2d4a7 100644 --- a/clang/include/clang/AST/TextNodeDumper.h +++ b/clang/include/clang/AST/TextNodeDumper.h @@ -425,6 +425,7 @@ public: void VisitOpenACCCacheConstruct(const OpenACCCacheConstruct *S); void VisitOpenACCAsteriskSizeExpr(const OpenACCAsteriskSizeExpr *S); void VisitOpenACCDeclareDecl(const OpenACCDeclareDecl *D); + void VisitOpenACCRoutineDecl(const OpenACCRoutineDecl *D); 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 458747a..dc9b462 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -5026,3 +5026,11 @@ def Atomic : StmtAttr { let Documentation = [AtomicDocs]; let StrictEnumParameters = 1; } + +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. + let Spellings = []; + let Subjects = SubjectList<[Function]>; + let Documentation = [InternalOnly]; +} diff --git a/clang/include/clang/Basic/DeclNodes.td b/clang/include/clang/Basic/DeclNodes.td index 227184d..20debd67 100644 --- a/clang/include/clang/Basic/DeclNodes.td +++ b/clang/include/clang/Basic/DeclNodes.td @@ -112,3 +112,4 @@ def RequiresExprBody : DeclNode<Decl>, DeclContext; def LifetimeExtendedTemporary : DeclNode<Decl>; def HLSLBuffer : DeclNode<Named, "HLSLBuffer">, DeclContext; def OpenACCDeclare : DeclNode<Decl, "#pragma acc declare">; +def OpenACCRoutine : DeclNode<Decl, "#pragma acc routine">; diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 5e5902cd..d712de6 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12821,8 +12821,9 @@ def err_wasm_builtin_arg_must_be_integer_type : Error < "%ordinal0 argument must be an integer">; // OpenACC diagnostics. -def warn_acc_construct_unimplemented - : Warning<"OpenACC construct '%0' not yet implemented, pragma ignored">, +def warn_acc_routine_unimplemented + : Warning<"OpenACC construct 'routine' with implicit function not yet " + "implemented, pragma ignored">, InGroup<SourceUsesOpenACC>; def warn_acc_clause_unimplemented : Warning<"OpenACC clause '%0' not yet implemented, clause ignored">, @@ -13060,6 +13061,13 @@ def err_acc_declare_same_scope def err_acc_multiple_references : Error<"variable referenced in '%0' clause of OpenACC 'declare' directive " "was already referenced">; +def err_acc_routine_not_func + : Error<"OpenACC routine name '%0' does not name a function">; +def err_acc_routine_overload_set + : Error<"OpenACC routine name '%0' names a set of overloads">; +def err_acc_magic_static_in_routine + : Error<"function static variables are not permitted in functions to which " + "an OpenACC 'routine' directive applies">; // AMDGCN builtins diagnostics def err_amdgcn_global_load_lds_size_invalid_value : Error<"invalid size value">; diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 80177996b..19f35f7 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -13436,6 +13436,10 @@ public: bool ForCallExpr = false); ExprResult SubstExpr(Expr *E, const MultiLevelTemplateArgumentList &TemplateArgs); + /// Substitute an expression as if it is a address-of-operand, which makes it + /// act like a CXXIdExpression rather than an attempt to call. + ExprResult SubstCXXIdExpr(Expr *E, + const MultiLevelTemplateArgumentList &TemplateArgs); // A RAII type used by the TemplateDeclInstantiator and TemplateInstantiator // to disable constraint evaluation, then restore the state. diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h index 6edc0d6..442e53f 100644 --- a/clang/include/clang/Sema/SemaOpenACC.h +++ b/clang/include/clang/Sema/SemaOpenACC.h @@ -167,6 +167,11 @@ private: // contexts, we can just store the declaration and location of the reference. llvm::DenseMap<const clang::DeclaratorDecl *, SourceLocation> DeclareVarReferences; + // The 'routine' construct disallows magic-statics in a function referred to + // by a 'routine' directive. So record any of these that we see so we can + // check them later. + llvm::SmallDenseMap<const clang::FunctionDecl *, SourceLocation> + MagicStaticLocs; public: ComputeConstructInfo &getActiveComputeConstructInfo() { @@ -723,6 +728,7 @@ public: /// MiscLoc: First misc location, if necessary (not all constructs). /// Exprs: List of expressions on the construct itself, if necessary (not all /// constructs). + /// FuncRef: used only for Routine, this is the function being referenced. /// AK: The atomic kind of the directive, if necessary (atomic only) /// RParenLoc: Location of the right paren, if it exists (not on all /// constructs). @@ -735,23 +741,12 @@ public: OpenACCAtomicKind AK, SourceLocation RParenLoc, SourceLocation EndLoc, ArrayRef<OpenACCClause *> Clauses, StmtResult AssocStmt); - StmtResult ActOnEndStmtDirective( - OpenACCDirectiveKind K, SourceLocation StartLoc, SourceLocation DirLoc, - SourceLocation LParenLoc, SourceLocation MiscLoc, ArrayRef<Expr *> Exprs, - SourceLocation RParenLoc, SourceLocation EndLoc, - ArrayRef<OpenACCClause *> Clauses, StmtResult AssocStmt) { - return ActOnEndStmtDirective(K, StartLoc, DirLoc, LParenLoc, MiscLoc, Exprs, - OpenACCAtomicKind::None, RParenLoc, EndLoc, - Clauses, AssocStmt); - } - /// Called after the directive has been completely parsed, including the /// declaration group or associated statement. - DeclGroupRef ActOnEndDeclDirective(OpenACCDirectiveKind K, - SourceLocation StartLoc, - SourceLocation DirLoc, - SourceLocation EndLoc, - ArrayRef<OpenACCClause *> Clauses); + DeclGroupRef ActOnEndDeclDirective( + OpenACCDirectiveKind K, SourceLocation StartLoc, SourceLocation DirLoc, + SourceLocation LParenLoc, Expr *FuncRef, SourceLocation RParenLoc, + SourceLocation EndLoc, ArrayRef<OpenACCClause *> Clauses); /// Called when encountering an 'int-expr' for OpenACC, and manages /// conversions and diagnostics to 'int'. @@ -764,6 +759,9 @@ public: Expr *VarExpr); /// 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. + void ActOnVariableDeclarator(VarDecl *VD); // Called after 'ActOnVar' specifically for a 'link' clause, which has to do // some minor additional checks. @@ -773,6 +771,8 @@ public: // checked during both phases of template translation. bool CheckDeclareClause(SemaOpenACC::OpenACCParsedClause &Clause); + ExprResult ActOnRoutineName(Expr *RoutineName); + /// Called while semantically analyzing the reduction clause, ensuring the var /// is the correct kind of reference. ExprResult CheckReductionVar(OpenACCDirectiveKind DirectiveKind, diff --git a/clang/include/clang/Serialization/ASTBitCodes.h b/clang/include/clang/Serialization/ASTBitCodes.h index efb52cf..5cb9998 100644 --- a/clang/include/clang/Serialization/ASTBitCodes.h +++ b/clang/include/clang/Serialization/ASTBitCodes.h @@ -1531,7 +1531,10 @@ enum DeclCode { // An OpenACCDeclareDecl record. DECL_OPENACC_DECLARE, - DECL_LAST = DECL_OPENACC_DECLARE + // An OpenACCRoutineDecl record. + DECL_OPENACC_ROUTINE, + + DECL_LAST = DECL_OPENACC_ROUTINE }; /// Record codes for each kind of statement or expression. diff --git a/clang/lib/AST/DeclBase.cpp b/clang/lib/AST/DeclBase.cpp index 9d8eb07..6260b92 100644 --- a/clang/lib/AST/DeclBase.cpp +++ b/clang/lib/AST/DeclBase.cpp @@ -994,6 +994,7 @@ unsigned Decl::getIdentifierNamespaceForKind(Kind DeclKind) { case RequiresExprBody: case ImplicitConceptSpecialization: case OpenACCDeclare: + case OpenACCRoutine: // Never looked up by name. return 0; } diff --git a/clang/lib/AST/DeclOpenACC.cpp b/clang/lib/AST/DeclOpenACC.cpp index 568fd9d..f1e6770 100644 --- a/clang/lib/AST/DeclOpenACC.cpp +++ b/clang/lib/AST/DeclOpenACC.cpp @@ -31,3 +31,22 @@ OpenACCDeclareDecl::CreateDeserialized(ASTContext &Ctx, GlobalDeclID ID, return new (Ctx, ID, additionalSizeToAlloc<const OpenACCClause *>(NumClauses)) OpenACCDeclareDecl(NumClauses); } + +OpenACCRoutineDecl * +OpenACCRoutineDecl::Create(ASTContext &Ctx, DeclContext *DC, + SourceLocation StartLoc, SourceLocation DirLoc, + SourceLocation LParenLoc, Expr *FuncRef, + SourceLocation RParenLoc, SourceLocation EndLoc, + ArrayRef<const OpenACCClause *> Clauses) { + return new (Ctx, DC, + additionalSizeToAlloc<const OpenACCClause *>(Clauses.size())) + OpenACCRoutineDecl(DC, StartLoc, DirLoc, LParenLoc, FuncRef, RParenLoc, + EndLoc, Clauses); +} + +OpenACCRoutineDecl * +OpenACCRoutineDecl::CreateDeserialized(ASTContext &Ctx, GlobalDeclID ID, + unsigned NumClauses) { + return new (Ctx, ID, additionalSizeToAlloc<const OpenACCClause *>(NumClauses)) + OpenACCRoutineDecl(NumClauses); +} diff --git a/clang/lib/AST/DeclPrinter.cpp b/clang/lib/AST/DeclPrinter.cpp index 5c29956..959f7dc 100644 --- a/clang/lib/AST/DeclPrinter.cpp +++ b/clang/lib/AST/DeclPrinter.cpp @@ -113,6 +113,7 @@ namespace { void VisitHLSLBufferDecl(HLSLBufferDecl *D); void VisitOpenACCDeclareDecl(OpenACCDeclareDecl *D); + void VisitOpenACCRoutineDecl(OpenACCRoutineDecl *D); void printTemplateParameters(const TemplateParameterList *Params, bool OmitTemplateKW = false); @@ -497,7 +498,7 @@ void DeclPrinter::VisitDeclContext(DeclContext *DC, bool Indent) { isa<OMPDeclareMapperDecl>(*D) || isa<OMPRequiresDecl>(*D) || isa<OMPAllocateDecl>(*D)) Terminator = nullptr; - else if (isa<OpenACCDeclareDecl>(*D)) + else if (isa<OpenACCDeclareDecl, OpenACCRoutineDecl>(*D)) Terminator = nullptr; else if (isa<ObjCMethodDecl>(*D) && cast<ObjCMethodDecl>(*D)->hasBody()) Terminator = nullptr; @@ -1922,3 +1923,25 @@ void DeclPrinter::VisitOpenACCDeclareDecl(OpenACCDeclareDecl *D) { Printer.VisitClauseList(D->clauses()); } } +void DeclPrinter::VisitOpenACCRoutineDecl(OpenACCRoutineDecl *D) { + if (!D->isInvalidDecl()) { + Out << "#pragma acc routine"; + + if (D->hasNameSpecified()) { + 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>"; + + Out << ")"; + } + + OpenACCClausePrinter Printer(Out, Policy); + Printer.VisitClauseList(D->clauses()); + } +} diff --git a/clang/lib/AST/JSONNodeDumper.cpp b/clang/lib/AST/JSONNodeDumper.cpp index d692033..e5e7bd3 100644 --- a/clang/lib/AST/JSONNodeDumper.cpp +++ b/clang/lib/AST/JSONNodeDumper.cpp @@ -1363,6 +1363,7 @@ void JSONNodeDumper::VisitOpenACCAsteriskSizeExpr( const OpenACCAsteriskSizeExpr *E) {} void JSONNodeDumper::VisitOpenACCDeclareDecl(const OpenACCDeclareDecl *D) {} +void JSONNodeDumper::VisitOpenACCRoutineDecl(const OpenACCRoutineDecl *D) {} void JSONNodeDumper::VisitPredefinedExpr(const PredefinedExpr *PE) { JOS.attribute("name", PredefinedExpr::getIdentKindName(PE->getIdentKind())); diff --git a/clang/lib/AST/StmtPrinter.cpp b/clang/lib/AST/StmtPrinter.cpp index 263c1ba..facdc41 100644 --- a/clang/lib/AST/StmtPrinter.cpp +++ b/clang/lib/AST/StmtPrinter.cpp @@ -264,7 +264,8 @@ void StmtPrinter::VisitDeclStmt(DeclStmt *Node) { Indent(); PrintRawDeclStmt(Node); // Certain pragma declarations shouldn't have a semi-colon after them. - if (!Node->isSingleDecl() || !isa<OpenACCDeclareDecl>(Node->getSingleDecl())) + if (!Node->isSingleDecl() || + !isa<OpenACCDeclareDecl, OpenACCRoutineDecl>(Node->getSingleDecl())) OS << ";"; OS << NL; } diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp index 9b2abce..c07919c 100644 --- a/clang/lib/AST/TextNodeDumper.cpp +++ b/clang/lib/AST/TextNodeDumper.cpp @@ -3079,6 +3079,23 @@ void TextNodeDumper::VisitOpenACCDeclareDecl(const OpenACCDeclareDecl *D) { AddChild([=] { Visit(S); }); }); } +void TextNodeDumper::VisitOpenACCRoutineDecl(const OpenACCRoutineDecl *D) { + OS << " " << D->getDirectiveKind(); + + if (D->hasNameSpecified()) { + OS << " name_specified"; + dumpSourceRange(SourceRange{D->getLParenLoc(), D->getRParenLoc()}); + } + + AddChild([=] { Visit(D->getFunctionReference()); }); + + for (const OpenACCClause *C : D->clauses()) + AddChild([=] { + Visit(C); + for (const Stmt *S : C->children()) + AddChild([=] { Visit(S); }); + }); +} void TextNodeDumper::VisitEmbedExpr(const EmbedExpr *S) { AddChild("begin", [=] { OS << S->getStartingElementPos(); }); diff --git a/clang/lib/CodeGen/CGDecl.cpp b/clang/lib/CodeGen/CGDecl.cpp index 9cd5885..3ad9ebf 100644 --- a/clang/lib/CodeGen/CGDecl.cpp +++ b/clang/lib/CodeGen/CGDecl.cpp @@ -180,6 +180,8 @@ void CodeGenFunction::EmitDecl(const Decl &D) { case Decl::OpenACCDeclare: return CGM.EmitOpenACCDeclare(cast<OpenACCDeclareDecl>(&D), this); + case Decl::OpenACCRoutine: + return CGM.EmitOpenACCRoutine(cast<OpenACCRoutineDecl>(&D), this); case Decl::Typedef: // typedef int X; case Decl::TypeAlias: { // using X = int; [C++0x] @@ -2852,6 +2854,11 @@ void CodeGenModule::EmitOpenACCDeclare(const OpenACCDeclareDecl *D, // This is a no-op, we cna just ignore these declarations. } +void CodeGenModule::EmitOpenACCRoutine(const OpenACCRoutineDecl *D, + CodeGenFunction *CGF) { + // This is a no-op, we cna just ignore these declarations. +} + void CodeGenModule::EmitOMPRequiresDecl(const OMPRequiresDecl *D) { getOpenMPRuntime().processRequiresDirective(D); } diff --git a/clang/lib/CodeGen/CodeGenModule.h b/clang/lib/CodeGen/CodeGenModule.h index a72b4f7..83bb5bc 100644 --- a/clang/lib/CodeGen/CodeGenModule.h +++ b/clang/lib/CodeGen/CodeGenModule.h @@ -1575,6 +1575,8 @@ public: // Emit code for the OpenACC Declare declaration. void EmitOpenACCDeclare(const OpenACCDeclareDecl *D, CodeGenFunction *CGF); + // Emit code for the OpenACC Routine declaration. + void EmitOpenACCRoutine(const OpenACCRoutineDecl *D, CodeGenFunction *CGF); /// Emit a code for requires directive. /// \param D Requires declaration diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp index 1417caa..754bd26 100644 --- a/clang/lib/Parse/ParseOpenACC.cpp +++ b/clang/lib/Parse/ParseOpenACC.cpp @@ -1441,6 +1441,7 @@ Parser::ParseOpenACCDirective() { Parser::OpenACCWaitParseInfo WaitInfo; Parser::OpenACCCacheParseInfo CacheInfo; OpenACCAtomicKind AtomicKind = OpenACCAtomicKind::None; + ExprResult RoutineName; getActions().OpenACC().ActOnConstruct(DirKind, DirLoc); @@ -1464,13 +1465,16 @@ Parser::ParseOpenACCDirective() { case OpenACCDirectiveKind::Routine: { // Routine has an optional paren-wrapped name of a function in the local // scope. We parse the name, emitting any diagnostics - ExprResult RoutineName = ParseOpenACCIDExpression(); + RoutineName = ParseOpenACCIDExpression(); // If the routine name is invalid, just skip until the closing paren to // recover more gracefully. - if (RoutineName.isInvalid()) + if (!RoutineName.isUsable()) { T.skipToEnd(); - else + } else { T.consumeClose(); + RoutineName = + getActions().OpenACC().ActOnRoutineName(RoutineName.get()); + } break; } case OpenACCDirectiveKind::Cache: @@ -1496,19 +1500,27 @@ Parser::ParseOpenACCDirective() { } // Parses the list of clauses, if present, plus set up return value. - OpenACCDirectiveParseInfo ParseInfo{ - DirKind, - StartLoc, - DirLoc, - T.getOpenLocation(), - T.getCloseLocation(), - /*EndLoc=*/SourceLocation{}, - (DirKind == OpenACCDirectiveKind::Wait ? WaitInfo.QueuesLoc - : CacheInfo.ReadOnlyLoc), - AtomicKind, - (DirKind == OpenACCDirectiveKind::Wait ? WaitInfo.getAllExprs() - : CacheInfo.Vars), - ParseOpenACCClauseList(DirKind)}; + OpenACCDirectiveParseInfo ParseInfo{DirKind, + StartLoc, + DirLoc, + T.getOpenLocation(), + T.getCloseLocation(), + /*EndLoc=*/SourceLocation{}, + (DirKind == OpenACCDirectiveKind::Wait + ? WaitInfo.QueuesLoc + : CacheInfo.ReadOnlyLoc), + AtomicKind, + {}, + {}}; + + if (DirKind == OpenACCDirectiveKind::Wait) + ParseInfo.Exprs = WaitInfo.getAllExprs(); + else if (DirKind == OpenACCDirectiveKind::Cache) + ParseInfo.Exprs = std::move(CacheInfo.Vars); + else if (DirKind == OpenACCDirectiveKind::Routine && RoutineName.isUsable()) + ParseInfo.Exprs = llvm::SmallVector<Expr *>(1, RoutineName.get()); + + ParseInfo.Clauses = ParseOpenACCClauseList(DirKind); assert(Tok.is(tok::annot_pragma_openacc_end) && "Didn't parse all OpenACC Clauses"); @@ -1532,8 +1544,9 @@ Parser::DeclGroupPtrTy Parser::ParseOpenACCDirectiveDecl() { return nullptr; return DeclGroupPtrTy::make(getActions().OpenACC().ActOnEndDeclDirective( - DirInfo.DirKind, DirInfo.StartLoc, DirInfo.DirLoc, DirInfo.EndLoc, - DirInfo.Clauses)); + DirInfo.DirKind, DirInfo.StartLoc, DirInfo.DirLoc, DirInfo.LParenLoc, + DirInfo.Exprs.empty() ? nullptr : DirInfo.Exprs.front(), + DirInfo.RParenLoc, DirInfo.EndLoc, DirInfo.Clauses)); } // Parse OpenACC Directive on a Statement. diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 36de02d..8c5125b8 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -50,6 +50,7 @@ #include "clang/Sema/SemaHLSL.h" #include "clang/Sema/SemaInternal.h" #include "clang/Sema/SemaObjC.h" +#include "clang/Sema/SemaOpenACC.h" #include "clang/Sema/SemaOpenMP.h" #include "clang/Sema/SemaPPC.h" #include "clang/Sema/SemaRISCV.h" @@ -7977,6 +7978,8 @@ NamedDecl *Sema::ActOnVariableDeclarator( if (getLangOpts().HLSL) HLSL().ActOnVariableDeclarator(NewVD); + if (getLangOpts().OpenACC) + OpenACC().ActOnVariableDeclarator(NewVD); // FIXME: This is probably the wrong location to be doing this and we should // probably be doing this for more attributes (especially for function diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp index 85fa631..22b386c 100644 --- a/clang/lib/Sema/SemaOpenACC.cpp +++ b/clang/lib/Sema/SemaOpenACC.cpp @@ -314,42 +314,8 @@ void SemaOpenACC::ActOnConstruct(OpenACCDirectiveKind K, SemaRef.PushExpressionEvaluationContext( Sema::ExpressionEvaluationContext::PotentiallyEvaluated); - switch (K) { - case OpenACCDirectiveKind::Invalid: - // Nothing to do here, an invalid kind has nothing we can check here. We - // want to continue parsing clauses as far as we can, so we will just - // ensure that we can still work and don't check any construct-specific - // rules anywhere. - break; - case OpenACCDirectiveKind::Parallel: - case OpenACCDirectiveKind::Serial: - case OpenACCDirectiveKind::Kernels: - case OpenACCDirectiveKind::ParallelLoop: - case OpenACCDirectiveKind::SerialLoop: - case OpenACCDirectiveKind::KernelsLoop: - case OpenACCDirectiveKind::Loop: - case OpenACCDirectiveKind::Data: - case OpenACCDirectiveKind::EnterData: - case OpenACCDirectiveKind::ExitData: - case OpenACCDirectiveKind::HostData: - case OpenACCDirectiveKind::Init: - case OpenACCDirectiveKind::Shutdown: - case OpenACCDirectiveKind::Set: - case OpenACCDirectiveKind::Update: - case OpenACCDirectiveKind::Cache: - case OpenACCDirectiveKind::Atomic: - case OpenACCDirectiveKind::Declare: - // Nothing to do here, there is no real legalization that needs to happen - // here as these constructs do not take any arguments. - break; - case OpenACCDirectiveKind::Wait: - // Nothing really to do here, the arguments to the 'wait' should have - // already been handled by the time we get here. - break; - default: - Diag(DirLoc, diag::warn_acc_construct_unimplemented) << K; - break; - } + // There is nothing do do here as all we have at this point is the name of the + // construct itself. } ExprResult SemaOpenACC::ActOnIntExpr(OpenACCDirectiveKind DK, @@ -1480,8 +1446,108 @@ std::string GetListOfClauses(llvm::ArrayRef<OpenACCClauseKind> Clauses) { OS << " or \'" << Clauses.back() << '\''; return Output; } + +// Helper that should mirror ActOnRoutineName to get the FunctionDecl out for +// magic-static checking. +FunctionDecl *getFunctionFromRoutineName(Expr *RoutineName) { + if (!RoutineName) + return nullptr; + RoutineName = RoutineName->IgnoreParenImpCasts(); + if (isa<RecoveryExpr>(RoutineName)) { + return nullptr; + } else if (isa<DependentScopeDeclRefExpr, CXXDependentScopeMemberExpr>( + RoutineName)) { + return nullptr; + } else if (auto *DRE = dyn_cast<DeclRefExpr>(RoutineName)) { + ValueDecl *VD = DRE->getDecl(); + + if (auto *FD = dyn_cast<FunctionDecl>(VD)) + return FD; + + // Allow lambdas. + if (auto *VarD = dyn_cast<VarDecl>(VD)) { + if (auto *RD = VarD->getType()->getAsCXXRecordDecl()) { + if (RD->isGenericLambda()) + return nullptr; + if (RD->isLambda()) + return RD->getLambdaCallOperator(); + } + } + return nullptr; + } else if (isa<OverloadExpr>(RoutineName)) { + return nullptr; + } + return nullptr; +} + } // namespace +ExprResult SemaOpenACC::ActOnRoutineName(Expr *RoutineName) { + assert(RoutineName && "Routine name cannot be null here"); + RoutineName = RoutineName->IgnoreParenImpCasts(); + + if (isa<RecoveryExpr>(RoutineName)) { + // This has already been diagnosed, so we can skip it. + return ExprError(); + } else if (isa<DependentScopeDeclRefExpr, CXXDependentScopeMemberExpr>( + RoutineName)) { + // These are dependent and we can't really check them, so delay until + // instantiation. + return RoutineName; + } else if (const auto *DRE = dyn_cast<DeclRefExpr>(RoutineName)) { + const ValueDecl *VD = DRE->getDecl(); + + if (const auto *FD = dyn_cast<FunctionDecl>(VD)) + return 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()) + 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; + } + } + + Diag(RoutineName->getBeginLoc(), diag::err_acc_routine_not_func) + << RoutineName; + return ExprError(); + } else if (isa<OverloadExpr>(RoutineName)) { + // This happens in function templates, even when the template arguments are + // fully specified. We could possibly do some sort of matching to make sure + // that this is looked up/deduced, but GCC does not do this, so there + // doesn't seem to be a good reason for us to do it either. + Diag(RoutineName->getBeginLoc(), diag::err_acc_routine_overload_set) + << RoutineName; + return ExprError(); + } + + Diag(RoutineName->getBeginLoc(), diag::err_acc_routine_not_func) + << RoutineName; + return ExprError(); +} +void SemaOpenACC::ActOnVariableDeclarator(VarDecl *VD) { + if (!VD->isStaticLocal()) + return; + + if (const auto *FD = dyn_cast<FunctionDecl>(getCurContext())) { + if (const auto *A = FD->getAttr<OpenACCRoutineAnnotAttr>()) { + Diag(VD->getBeginLoc(), diag::err_acc_magic_static_in_routine); + Diag(A->getLocation(), diag::note_acc_construct_here) + << OpenACCDirectiveKind::Routine; + } + MagicStaticLocs.insert({FD, VD->getBeginLoc()}); + } +} + bool SemaOpenACC::ActOnStartStmtDirective( OpenACCDirectiveKind K, SourceLocation StartLoc, ArrayRef<const OpenACCClause *> Clauses) { @@ -1607,8 +1673,6 @@ StmtResult SemaOpenACC::ActOnEndStmtDirective( SourceLocation EndLoc, ArrayRef<OpenACCClause *> Clauses, StmtResult AssocStmt) { switch (K) { - default: - return StmtEmpty(); case OpenACCDirectiveKind::Invalid: return StmtError(); case OpenACCDirectiveKind::Parallel: @@ -1681,12 +1745,14 @@ StmtResult SemaOpenACC::ActOnEndStmtDirective( LParenLoc, MiscLoc, Exprs, RParenLoc, EndLoc); } + case OpenACCDirectiveKind::Routine: case OpenACCDirectiveKind::Declare: { - // Declare is a declaration directive, but can be used here as long as we - // wrap it in a DeclStmt. So make sure we do that here. + // 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, EndLoc, Clauses); return SemaRef.ActOnDeclStmt(DeclGroupPtrTy::make(DR), StartLoc, EndLoc); } } @@ -1773,11 +1839,13 @@ bool SemaOpenACC::ActOnStartDeclDirective(OpenACCDirectiveKind K, // clause arguments or on any side effects of the evaluations. SemaRef.DiscardCleanupsInEvaluationContext(); SemaRef.PopExpressionEvaluationContext(); + return diagnoseConstructAppertainment(*this, K, StartLoc, /*IsStmt=*/false); } DeclGroupRef SemaOpenACC::ActOnEndDeclDirective( OpenACCDirectiveKind K, SourceLocation StartLoc, SourceLocation DirLoc, + SourceLocation LParenLoc, Expr *FuncRef, SourceLocation RParenLoc, SourceLocation EndLoc, ArrayRef<OpenACCClause *> Clauses) { switch (K) { default: @@ -1799,8 +1867,34 @@ 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{}; + } + + 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: + // 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; + } + FD->addAttr(OpenACCRoutineAnnotAttr::Create(getASTContext(), DirLoc)); + } + return DeclGroupRef{RoutineDecl}; + } + } llvm_unreachable("unhandled case in directive handling?"); } diff --git a/clang/lib/Sema/SemaTemplateInstantiate.cpp b/clang/lib/Sema/SemaTemplateInstantiate.cpp index fcb7671..19c27a7 100644 --- a/clang/lib/Sema/SemaTemplateInstantiate.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiate.cpp @@ -4499,6 +4499,17 @@ Sema::SubstExpr(Expr *E, const MultiLevelTemplateArgumentList &TemplateArgs) { } ExprResult +Sema::SubstCXXIdExpr(Expr *E, + const MultiLevelTemplateArgumentList &TemplateArgs) { + if (!E) + return E; + + TemplateInstantiator Instantiator(*this, TemplateArgs, SourceLocation(), + DeclarationName()); + return Instantiator.TransformAddressOfOperand(E); +} + +ExprResult Sema::SubstConstraintExpr(Expr *E, const MultiLevelTemplateArgumentList &TemplateArgs) { // FIXME: should call SubstExpr directly if this function is equivalent or diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp index 82e049c..7294e42 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -1218,8 +1218,38 @@ Decl *TemplateDeclInstantiator::VisitOpenACCDeclareDecl(OpenACCDeclareDecl *D) { return nullptr; DeclGroupRef Res = SemaRef.OpenACC().ActOnEndDeclDirective( + D->getDirectiveKind(), D->getBeginLoc(), D->getDirectiveLoc(), {}, + nullptr, {}, D->getEndLoc(), TransformedClauses); + + if (Res.isNull()) + return nullptr; + + return Res.getSingleDecl(); +} + +Decl *TemplateDeclInstantiator::VisitOpenACCRoutineDecl(OpenACCRoutineDecl *D) { + SemaRef.OpenACC().ActOnConstruct(D->getDirectiveKind(), D->getBeginLoc()); + llvm::SmallVector<OpenACCClause *> TransformedClauses = + InstantiateOpenACCClauseList(SemaRef, TemplateArgs, D->getDirectiveKind(), + D->clauses()); + + ExprResult FuncRef; + if (D->getFunctionReference()) { + FuncRef = SemaRef.SubstCXXIdExpr(D->getFunctionReference(), TemplateArgs); + if (FuncRef.isUsable()) + FuncRef = SemaRef.OpenACC().ActOnRoutineName(FuncRef.get()); + // We don't return early here, we leave the construct in the AST, even if + // the function decl is empty. + } + + if (SemaRef.OpenACC().ActOnStartDeclDirective(D->getDirectiveKind(), + D->getBeginLoc())) + return nullptr; + + DeclGroupRef Res = SemaRef.OpenACC().ActOnEndDeclDirective( D->getDirectiveKind(), D->getBeginLoc(), D->getDirectiveLoc(), - D->getEndLoc(), TransformedClauses); + D->getLParenLoc(), FuncRef.get(), D->getRParenLoc(), D->getEndLoc(), + TransformedClauses); if (Res.isNull()) return nullptr; diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h index cb2f335..b99acf6 100644 --- a/clang/lib/Sema/TreeTransform.h +++ b/clang/lib/Sema/TreeTransform.h @@ -4093,7 +4093,7 @@ public: StmtResult StrBlock) { return getSema().OpenACC().ActOnEndStmtDirective( K, BeginLoc, DirLoc, SourceLocation{}, SourceLocation{}, {}, - SourceLocation{}, EndLoc, Clauses, StrBlock); + OpenACCAtomicKind::None, SourceLocation{}, EndLoc, Clauses, StrBlock); } StmtResult RebuildOpenACCLoopConstruct(SourceLocation BeginLoc, @@ -4103,7 +4103,8 @@ public: StmtResult Loop) { return getSema().OpenACC().ActOnEndStmtDirective( OpenACCDirectiveKind::Loop, BeginLoc, DirLoc, SourceLocation{}, - SourceLocation{}, {}, SourceLocation{}, EndLoc, Clauses, Loop); + SourceLocation{}, {}, OpenACCAtomicKind::None, SourceLocation{}, EndLoc, + Clauses, Loop); } StmtResult RebuildOpenACCCombinedConstruct(OpenACCDirectiveKind K, @@ -4114,7 +4115,7 @@ public: StmtResult Loop) { return getSema().OpenACC().ActOnEndStmtDirective( K, BeginLoc, DirLoc, SourceLocation{}, SourceLocation{}, {}, - SourceLocation{}, EndLoc, Clauses, Loop); + OpenACCAtomicKind::None, SourceLocation{}, EndLoc, Clauses, Loop); } StmtResult RebuildOpenACCDataConstruct(SourceLocation BeginLoc, @@ -4124,7 +4125,8 @@ public: StmtResult StrBlock) { return getSema().OpenACC().ActOnEndStmtDirective( OpenACCDirectiveKind::Data, BeginLoc, DirLoc, SourceLocation{}, - SourceLocation{}, {}, SourceLocation{}, EndLoc, Clauses, StrBlock); + SourceLocation{}, {}, OpenACCAtomicKind::None, SourceLocation{}, EndLoc, + Clauses, StrBlock); } StmtResult @@ -4133,7 +4135,8 @@ public: ArrayRef<OpenACCClause *> Clauses) { return getSema().OpenACC().ActOnEndStmtDirective( OpenACCDirectiveKind::EnterData, BeginLoc, DirLoc, SourceLocation{}, - SourceLocation{}, {}, SourceLocation{}, EndLoc, Clauses, {}); + SourceLocation{}, {}, OpenACCAtomicKind::None, SourceLocation{}, EndLoc, + Clauses, {}); } StmtResult @@ -4142,7 +4145,8 @@ public: ArrayRef<OpenACCClause *> Clauses) { return getSema().OpenACC().ActOnEndStmtDirective( OpenACCDirectiveKind::ExitData, BeginLoc, DirLoc, SourceLocation{}, - SourceLocation{}, {}, SourceLocation{}, EndLoc, Clauses, {}); + SourceLocation{}, {}, OpenACCAtomicKind::None, SourceLocation{}, EndLoc, + Clauses, {}); } StmtResult RebuildOpenACCHostDataConstruct(SourceLocation BeginLoc, @@ -4152,7 +4156,8 @@ public: StmtResult StrBlock) { return getSema().OpenACC().ActOnEndStmtDirective( OpenACCDirectiveKind::HostData, BeginLoc, DirLoc, SourceLocation{}, - SourceLocation{}, {}, SourceLocation{}, EndLoc, Clauses, StrBlock); + SourceLocation{}, {}, OpenACCAtomicKind::None, SourceLocation{}, EndLoc, + Clauses, StrBlock); } StmtResult RebuildOpenACCInitConstruct(SourceLocation BeginLoc, @@ -4161,7 +4166,8 @@ public: ArrayRef<OpenACCClause *> Clauses) { return getSema().OpenACC().ActOnEndStmtDirective( OpenACCDirectiveKind::Init, BeginLoc, DirLoc, SourceLocation{}, - SourceLocation{}, {}, SourceLocation{}, EndLoc, Clauses, {}); + SourceLocation{}, {}, OpenACCAtomicKind::None, SourceLocation{}, EndLoc, + Clauses, {}); } StmtResult @@ -4170,7 +4176,8 @@ public: ArrayRef<OpenACCClause *> Clauses) { return getSema().OpenACC().ActOnEndStmtDirective( OpenACCDirectiveKind::Shutdown, BeginLoc, DirLoc, SourceLocation{}, - SourceLocation{}, {}, SourceLocation{}, EndLoc, Clauses, {}); + SourceLocation{}, {}, OpenACCAtomicKind::None, SourceLocation{}, EndLoc, + Clauses, {}); } StmtResult RebuildOpenACCSetConstruct(SourceLocation BeginLoc, @@ -4179,7 +4186,8 @@ public: ArrayRef<OpenACCClause *> Clauses) { return getSema().OpenACC().ActOnEndStmtDirective( OpenACCDirectiveKind::Set, BeginLoc, DirLoc, SourceLocation{}, - SourceLocation{}, {}, SourceLocation{}, EndLoc, Clauses, {}); + SourceLocation{}, {}, OpenACCAtomicKind::None, SourceLocation{}, EndLoc, + Clauses, {}); } StmtResult RebuildOpenACCUpdateConstruct(SourceLocation BeginLoc, @@ -4188,7 +4196,8 @@ public: ArrayRef<OpenACCClause *> Clauses) { return getSema().OpenACC().ActOnEndStmtDirective( OpenACCDirectiveKind::Update, BeginLoc, DirLoc, SourceLocation{}, - SourceLocation{}, {}, SourceLocation{}, EndLoc, Clauses, {}); + SourceLocation{}, {}, OpenACCAtomicKind::None, SourceLocation{}, EndLoc, + Clauses, {}); } StmtResult RebuildOpenACCWaitConstruct( @@ -4201,7 +4210,7 @@ public: Exprs.insert(Exprs.end(), QueueIdExprs.begin(), QueueIdExprs.end()); return getSema().OpenACC().ActOnEndStmtDirective( OpenACCDirectiveKind::Wait, BeginLoc, DirLoc, LParenLoc, QueuesLoc, - Exprs, RParenLoc, EndLoc, Clauses, {}); + Exprs, OpenACCAtomicKind::None, RParenLoc, EndLoc, Clauses, {}); } StmtResult RebuildOpenACCCacheConstruct( @@ -4210,7 +4219,7 @@ public: SourceLocation RParenLoc, SourceLocation EndLoc) { return getSema().OpenACC().ActOnEndStmtDirective( OpenACCDirectiveKind::Cache, BeginLoc, DirLoc, LParenLoc, ReadOnlyLoc, - VarList, RParenLoc, EndLoc, {}, {}); + VarList, OpenACCAtomicKind::None, RParenLoc, EndLoc, {}, {}); } StmtResult RebuildOpenACCAtomicConstruct(SourceLocation BeginLoc, diff --git a/clang/lib/Serialization/ASTCommon.cpp b/clang/lib/Serialization/ASTCommon.cpp index ae8d0a6..320ee0e 100644 --- a/clang/lib/Serialization/ASTCommon.cpp +++ b/clang/lib/Serialization/ASTCommon.cpp @@ -459,6 +459,7 @@ bool serialization::isRedeclarableDeclKind(unsigned Kind) { case Decl::UnresolvedUsingIfExists: case Decl::HLSLBuffer: case Decl::OpenACCDeclare: + case Decl::OpenACCRoutine: return false; // These indirectly derive from Redeclarable<T> but are not actually diff --git a/clang/lib/Serialization/ASTReaderDecl.cpp b/clang/lib/Serialization/ASTReaderDecl.cpp index 2a580c4..262cb0f 100644 --- a/clang/lib/Serialization/ASTReaderDecl.cpp +++ b/clang/lib/Serialization/ASTReaderDecl.cpp @@ -415,6 +415,7 @@ public: void VisitLifetimeExtendedTemporaryDecl(LifetimeExtendedTemporaryDecl *D); void VisitOpenACCDeclareDecl(OpenACCDeclareDecl *D); + void VisitOpenACCRoutineDecl(OpenACCRoutineDecl *D); void VisitDeclContext(DeclContext *DC, uint64_t &LexicalOffset, uint64_t &VisibleOffset, uint64_t &ModuleLocalOffset, @@ -3108,6 +3109,15 @@ void ASTDeclReader::VisitOpenACCDeclareDecl(OpenACCDeclareDecl *D) { D->EndLoc = Record.readSourceLocation(); Record.readOpenACCClauseList(D->Clauses); } +void ASTDeclReader::VisitOpenACCRoutineDecl(OpenACCRoutineDecl *D) { + VisitDecl(D); + D->DirKind = Record.readEnum<OpenACCDirectiveKind>(); + D->DirectiveLoc = Record.readSourceLocation(); + D->EndLoc = Record.readSourceLocation(); + D->ParensLoc = Record.readSourceRange(); + D->FuncRef = Record.readExpr(); + Record.readOpenACCClauseList(D->Clauses); +} //===----------------------------------------------------------------------===// // Attribute Reading @@ -4217,6 +4227,9 @@ Decl *ASTReader::ReadDeclRecord(GlobalDeclID ID) { case DECL_OPENACC_DECLARE: D = OpenACCDeclareDecl::CreateDeserialized(Context, ID, Record.readInt()); break; + case DECL_OPENACC_ROUTINE: + D = OpenACCRoutineDecl::CreateDeserialized(Context, ID, Record.readInt()); + break; } assert(D && "Unknown declaration reading AST file"); diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp index e514538..c3e67ff 100644 --- a/clang/lib/Serialization/ASTWriter.cpp +++ b/clang/lib/Serialization/ASTWriter.cpp @@ -1136,6 +1136,7 @@ void ASTWriter::WriteBlockInfoBlock() { RECORD(DECL_OMP_ALLOCATE); RECORD(DECL_HLSL_BUFFER); RECORD(DECL_OPENACC_DECLARE); + RECORD(DECL_OPENACC_ROUTINE); // Statements and Exprs can occur in the Decls and Types block. AddStmtsExprs(Stream, Record); diff --git a/clang/lib/Serialization/ASTWriterDecl.cpp b/clang/lib/Serialization/ASTWriterDecl.cpp index 88dd212..d455134 100644 --- a/clang/lib/Serialization/ASTWriterDecl.cpp +++ b/clang/lib/Serialization/ASTWriterDecl.cpp @@ -177,6 +177,7 @@ namespace clang { void VisitOMPCapturedExprDecl(OMPCapturedExprDecl *D); void VisitOpenACCDeclareDecl(OpenACCDeclareDecl *D); + void VisitOpenACCRoutineDecl(OpenACCRoutineDecl *D); /// Add an Objective-C type parameter list to the given record. void AddObjCTypeParamList(ObjCTypeParamList *typeParams) { @@ -2281,6 +2282,17 @@ void ASTDeclWriter::VisitOpenACCDeclareDecl(OpenACCDeclareDecl *D) { Record.writeOpenACCClauseList(D->clauses()); Code = serialization::DECL_OPENACC_DECLARE; } +void ASTDeclWriter::VisitOpenACCRoutineDecl(OpenACCRoutineDecl *D) { + Record.writeUInt32(D->clauses().size()); + VisitDecl(D); + Record.writeEnum(D->DirKind); + Record.AddSourceLocation(D->DirectiveLoc); + Record.AddSourceLocation(D->EndLoc); + Record.AddSourceRange(D->ParensLoc); + Record.AddStmt(D->FuncRef); + Record.writeOpenACCClauseList(D->clauses()); + Code = serialization::DECL_OPENACC_ROUTINE; +} //===----------------------------------------------------------------------===// // ASTWriter Implementation diff --git a/clang/test/AST/ast-print-openacc-routine-construct.cpp b/clang/test/AST/ast-print-openacc-routine-construct.cpp new file mode 100644 index 0000000..64628d5 --- /dev/null +++ b/clang/test/AST/ast-print-openacc-routine-construct.cpp @@ -0,0 +1,89 @@ +// RUN: %clang_cc1 -fopenacc -ast-print %s -o - | FileCheck %s + +auto Lambda = [](){}; +// CHECK: #pragma acc routine(Lambda) +#pragma acc routine(Lambda) +int function(); +// CHECK: #pragma acc routine(function) +#pragma acc routine (function) + +namespace NS { + int NSFunc(); +auto Lambda = [](){}; +} +// CHECK: #pragma acc routine(NS::NSFunc) +#pragma acc routine(NS::NSFunc) +// CHECK: #pragma acc routine(NS::Lambda) +#pragma acc routine(NS::Lambda) + +struct S { + void MemFunc(); + static void StaticMemFunc(); + constexpr static auto Lambda = [](){}; +// CHECK: #pragma acc routine(S::MemFunc) +#pragma acc routine(S::MemFunc) +// CHECK: #pragma acc routine(S::StaticMemFunc) +#pragma acc routine(S::StaticMemFunc) +// CHECK: #pragma acc routine(S::Lambda) +#pragma acc routine(S::Lambda) + +// CHECK: #pragma acc routine(MemFunc) +#pragma acc routine(MemFunc) +// CHECK: #pragma acc routine(StaticMemFunc) +#pragma acc routine(StaticMemFunc) +// CHECK: #pragma acc routine(Lambda) +#pragma acc routine(Lambda) +}; + +// CHECK: #pragma acc routine(S::MemFunc) +#pragma acc routine(S::MemFunc) +// CHECK: #pragma acc routine(S::StaticMemFunc) +#pragma acc routine(S::StaticMemFunc) +// CHECK: #pragma acc routine(S::Lambda) +#pragma acc routine(S::Lambda) + +template<typename T> +struct DepS { + void MemFunc(); + static void StaticMemFunc(); + constexpr static auto Lambda = [](){}; + +// CHECK: #pragma acc routine(Lambda) +#pragma acc routine(Lambda) +// CHECK: #pragma acc routine(MemFunc) +#pragma acc routine(MemFunc) +// CHECK: #pragma acc routine(StaticMemFunc) +#pragma acc routine(StaticMemFunc) + +// CHECK: #pragma acc routine(DepS<T>::Lambda) +#pragma acc routine(DepS::Lambda) +// CHECK: #pragma acc routine(DepS<T>::MemFunc) +#pragma acc routine(DepS::MemFunc) +// CHECK: #pragma acc routine(DepS<T>::StaticMemFunc) +#pragma acc routine(DepS::StaticMemFunc) + +// CHECK: #pragma acc routine(DepS<T>::Lambda) +#pragma acc routine(DepS<T>::Lambda) +// CHECK: #pragma acc routine(DepS<T>::MemFunc) +#pragma acc routine(DepS<T>::MemFunc) +// CHECK: #pragma acc routine(DepS<T>::StaticMemFunc) +#pragma acc routine(DepS<T>::StaticMemFunc) +}; + +// CHECK: #pragma acc routine(DepS<int>::Lambda) +#pragma acc routine(DepS<int>::Lambda) +// CHECK: #pragma acc routine(DepS<int>::MemFunc) +#pragma acc routine(DepS<int>::MemFunc) +// CHECK: #pragma acc routine(DepS<int>::StaticMemFunc) +#pragma acc routine(DepS<int>::StaticMemFunc) + + +template<typename T> +void TemplFunc() { +// CHECK: #pragma acc routine(T::MemFunc) +#pragma acc routine(T::MemFunc) +// CHECK: #pragma acc routine(T::StaticMemFunc) +#pragma acc routine(T::StaticMemFunc) +// CHECK: #pragma acc routine(T::Lambda) +#pragma acc routine(T::Lambda) +} diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c index 7c987ba..ad894c3 100644 --- a/clang/test/ParserOpenACC/parse-clauses.c +++ b/clang/test/ParserOpenACC/parse-clauses.c @@ -1265,38 +1265,34 @@ void Gang() { // expected-warning@+4{{OpenACC clause 'vector' not yet implemented, clause ignored}} // expected-warning@+3{{OpenACC clause 'seq' not yet implemented, clause ignored}} // expected-warning@+2{{OpenACC clause 'nohost' not yet implemented, clause ignored}} - // expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} + // expected-warning@+1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}} #pragma acc routine worker, vector, seq, nohost void bar(); - // expected-warning@+5{{OpenACC clause 'worker' not yet implemented, clause ignored}} - // expected-warning@+4{{OpenACC clause 'vector' not yet implemented, clause ignored}} - // expected-warning@+3{{OpenACC clause 'seq' not yet implemented, clause ignored}} - // expected-warning@+2{{OpenACC clause 'nohost' not yet implemented, clause ignored}} - // expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} + // expected-warning@+4{{OpenACC clause 'worker' not yet implemented, clause ignored}} + // expected-warning@+3{{OpenACC clause 'vector' not yet implemented, clause ignored}} + // expected-warning@+2{{OpenACC clause 'seq' not yet implemented, clause ignored}} + // expected-warning@+1{{OpenACC clause 'nohost' not yet implemented, clause ignored}} #pragma acc routine(bar) worker, vector, seq, nohost // Bind Clause Parsing. // expected-error@+2{{expected '('}} - // expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} + // expected-warning@+1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}} #pragma acc routine bind void BCP1(); - // expected-error@+2{{expected identifier or string literal}} - // expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} + // expected-error@+1{{expected identifier or string literal}} #pragma acc routine(BCP1) bind() // expected-warning@+2{{OpenACC clause 'bind' not yet implemented, clause ignored}} - // expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} + // expected-warning@+1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}} #pragma acc routine bind("ReductionClauseParsing") void BCP2(); - // expected-warning@+2{{OpenACC clause 'bind' not yet implemented, clause ignored}} - // expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} + // expected-warning@+1{{OpenACC clause 'bind' not yet implemented, clause ignored}} #pragma acc routine(BCP1) bind(BCP2) - // expected-error@+2{{use of undeclared identifier 'unknown_thing'}} - // expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} + // expected-error@+1{{use of undeclared identifier 'unknown_thing'}} #pragma acc routine(BCP1) bind(unknown_thing) diff --git a/clang/test/ParserOpenACC/parse-clauses.cpp b/clang/test/ParserOpenACC/parse-clauses.cpp index 770bc3b..ff747f3 100644 --- a/clang/test/ParserOpenACC/parse-clauses.cpp +++ b/clang/test/ParserOpenACC/parse-clauses.cpp @@ -65,48 +65,36 @@ void function(); } - // expected-warning@+2{{OpenACC clause 'bind' not yet implemented, clause ignored}} - // expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} + // expected-warning@+1{{OpenACC clause 'bind' not yet implemented, clause ignored}} #pragma acc routine(use) bind(NS::NSFunc) - // expected-error@+3{{'RecordTy' does not refer to a value}} + // expected-error@+2{{'RecordTy' does not refer to a value}} // expected-note@#RecTy{{declared here}} - // expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} #pragma acc routine(use) bind(NS::RecordTy) - // expected-error@+4{{'Value' is a private member of 'NS::RecordTy'}} + // expected-error@+3{{'Value' is a private member of 'NS::RecordTy'}} // expected-note@#VAL{{implicitly declared private here}} - // expected-warning@+2{{OpenACC clause 'bind' not yet implemented, clause ignored}} - // expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} + // expected-warning@+1{{OpenACC clause 'bind' not yet implemented, clause ignored}} #pragma acc routine(use) bind(NS::RecordTy::Value) - // expected-warning@+2{{OpenACC clause 'bind' not yet implemented, clause ignored}} - // expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} + // expected-warning@+1{{OpenACC clause 'bind' not yet implemented, clause ignored}} #pragma acc routine(use) bind(NS::RecordTy::ValuePub) - // expected-warning@+2{{OpenACC clause 'bind' not yet implemented, clause ignored}} - // expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} + // expected-warning@+1{{OpenACC clause 'bind' not yet implemented, clause ignored}} #pragma acc routine(use) bind(NS::TemplTy<int>) - // expected-error@+2{{no member named 'unknown' in namespace 'NS'}} - // expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} + // expected-error@+1{{no member named 'unknown' in namespace 'NS'}} #pragma acc routine(use) bind(NS::unknown<int>) - // expected-warning@+2{{OpenACC clause 'bind' not yet implemented, clause ignored}} - // expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} + // expected-warning@+1{{OpenACC clause 'bind' not yet implemented, clause ignored}} #pragma acc routine(use) bind(NS::function) - // expected-error@+4{{'priv_mem_function' is a private member of 'NS::RecordTy'}} + // expected-error@+3{{'priv_mem_function' is a private member of 'NS::RecordTy'}} // expected-note@#PrivMemFun{{implicitly declared private here}} - // expected-warning@+2{{OpenACC clause 'bind' not yet implemented, clause ignored}} - // expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} + // expected-warning@+1{{OpenACC clause 'bind' not yet implemented, clause ignored}} #pragma acc routine(use) bind(NS::RecordTy::priv_mem_function) - // expected-warning@+2{{OpenACC clause 'bind' not yet implemented, clause ignored}} - // expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} + // expected-warning@+1{{OpenACC clause 'bind' not yet implemented, clause ignored}} #pragma acc routine(use) bind(NS::RecordTy::mem_function) - // expected-error@+2{{string literal with user-defined suffix cannot be used here}} - // expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} + // expected-error@+1{{string literal with user-defined suffix cannot be used here}} #pragma acc routine(use) bind("unknown udl"_UDL) - // expected-warning@+3{{encoding prefix 'u' on an unevaluated string literal has no effect}} - // expected-warning@+2{{OpenACC clause 'bind' not yet implemented, clause ignored}} - // expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} + // expected-warning@+2{{encoding prefix 'u' on an unevaluated string literal has no effect}} + // expected-warning@+1{{OpenACC clause 'bind' not yet implemented, clause ignored}} #pragma acc routine(use) bind(u"16 bits") - // expected-warning@+3{{encoding prefix 'U' on an unevaluated string literal has no effect}} - // expected-warning@+2{{OpenACC clause 'bind' not yet implemented, clause ignored}} - // expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} + // expected-warning@+2{{encoding prefix 'U' on an unevaluated string literal has no effect}} + // expected-warning@+1{{OpenACC clause 'bind' not yet implemented, clause ignored}} #pragma acc routine(use) bind(U"32 bits") diff --git a/clang/test/ParserOpenACC/parse-constructs.c b/clang/test/ParserOpenACC/parse-constructs.c index 9762209..107d48a 100644 --- a/clang/test/ParserOpenACC/parse-constructs.c +++ b/clang/test/ParserOpenACC/parse-constructs.c @@ -150,37 +150,30 @@ void func() { for(;;){} } -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} +// expected-warning@+1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}} #pragma acc routine void routine_func(); // expected-error@+2{{invalid OpenACC clause 'clause'}} -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} +// expected-warning@+1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}} #pragma acc routine clause list void routine_func(); -// expected-error@+2{{use of undeclared identifier 'func_name'}} -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} +// expected-error@+1{{use of undeclared identifier 'func_name'}} #pragma acc routine (func_name) -// expected-error@+3{{use of undeclared identifier 'func_name'}} -// expected-error@+2{{invalid OpenACC clause 'clause'}} -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} +// expected-error@+2{{use of undeclared identifier 'func_name'}} +// expected-error@+1{{invalid OpenACC clause 'clause'}} #pragma acc routine (func_name) clause list -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} #pragma acc routine (routine_func) -// expected-error@+2{{invalid OpenACC clause 'clause'}} -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} +// expected-error@+1{{invalid OpenACC clause 'clause'}} #pragma acc routine (routine_func) clause list -// expected-error@+3{{expected ')'}} -// expected-note@+2{{to match this '('}} -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} +// expected-error@+2{{expected ')'}} +// expected-note@+1{{to match this '('}} #pragma acc routine (routine_func()) -// expected-error@+2{{expected identifier}} -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} +// expected-error@+1{{expected identifier}} #pragma acc routine() -// expected-error@+2{{expected identifier}} -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} +// expected-error@+1{{expected identifier}} #pragma acc routine(int) diff --git a/clang/test/ParserOpenACC/parse-constructs.cpp b/clang/test/ParserOpenACC/parse-constructs.cpp index 42fd09d..7892fb2 100644 --- a/clang/test/ParserOpenACC/parse-constructs.cpp +++ b/clang/test/ParserOpenACC/parse-constructs.cpp @@ -13,51 +13,41 @@ namespace NS { }; } -// expected-error@+2{{use of undeclared identifier 'foo'; did you mean 'NS::foo'?}} -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} +// expected-error@+1{{use of undeclared identifier 'foo'; did you mean 'NS::foo'?}} #pragma acc routine(foo) -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} #pragma acc routine(NS::foo) // expected-error@+2{{use of undeclared identifier 'templ'; did you mean 'NS::templ'?}} -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} +// expected-error@+1{{OpenACC routine name 'NS::templ' names a set of overloads}} #pragma acc routine(templ) -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} +// expected-error@+1{{OpenACC routine name 'NS::templ' names a set of overloads}} #pragma acc routine(NS::templ) // expected-error@+2{{use of undeclared identifier 'templ'; did you mean 'NS::templ'?}} -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} +// expected-error@+1{{OpenACC routine name 'NS::templ' names a set of overloads}} #pragma acc routine(templ<int>) -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} +// expected-error@+1{{OpenACC routine name 'NS::templ<int>' names a set of overloads}} #pragma acc routine(NS::templ<int>) -// expected-error@+2{{use of undeclared identifier 'T'}} -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} +// expected-error@+1{{use of undeclared identifier 'T'}} #pragma acc routine(templ<T>) -// expected-error@+2{{use of undeclared identifier 'T'}} -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} +// expected-error@+1{{use of undeclared identifier 'T'}} #pragma acc routine(NS::templ<T>) -// expected-error@+3{{expected ')'}} -// expected-note@+2{{to match this '('}} -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} +// expected-error@+2{{expected ')'}} +// expected-note@+1{{to match this '('}} #pragma acc routine (NS::foo()) -// expected-error@+2 {{expected unqualified-id}} -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} +// expected-error@+1 {{expected unqualified-id}} #pragma acc routine() -// expected-error@+2 {{expected unqualified-id}} -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} +// expected-error@+1 {{expected unqualified-id}} #pragma acc routine(int) -// expected-error@+3{{'C' does not refer to a value}} +// expected-error@+2{{'C' does not refer to a value}} // expected-note@#CDef{{declared here}} -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} #pragma acc routine (NS::C) -// expected-error@+3{{'private_mem_func' is a private member of 'NS::C'}} +// expected-error@+2{{'private_mem_func' is a private member of 'NS::C'}} // expected-note@#PrivateMemFunc{{implicitly declared private here}} -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} #pragma acc routine (NS::C::private_mem_func) -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} #pragma acc routine (NS::C::public_mem_func) diff --git a/clang/test/SemaOpenACC/routine-construct-ast.cpp b/clang/test/SemaOpenACC/routine-construct-ast.cpp new file mode 100644 index 0000000..1ef54c9 --- /dev/null +++ b/clang/test/SemaOpenACC/routine-construct-ast.cpp @@ -0,0 +1,204 @@ +// RUN: %clang_cc1 %s -fopenacc -ast-dump | FileCheck %s + +// Test this with PCH. +// RUN: %clang_cc1 %s -fopenacc -emit-pch -o %t %s +// RUN: %clang_cc1 %s -fopenacc -include-pch %t -ast-dump-all | FileCheck %s + +#ifndef PCH_HELPER +#define PCH_HELPER +auto Lambda = [](){}; +#pragma acc routine(Lambda) +// CHECK: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' '(lambda at +int function(); +#pragma acc routine (function) +// CHECK: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'function' 'int ()' + +namespace NS { + // CHECK-NEXT: NamespaceDecl + int NSFunc(); +auto Lambda = [](){}; +} +#pragma acc routine(NS::NSFunc) +// CHECK: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'NSFunc' 'int ()' +// CHECK-NEXT: NestedNameSpecifier{{.*}} 'NS' +#pragma acc routine(NS::Lambda) +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'NS::(lambda at +// CHECK-NEXT: NestedNameSpecifier{{.*}} 'NS' + +struct S { + void MemFunc(); + static void StaticMemFunc(); + constexpr static auto Lambda = [](){}; +#pragma acc routine(S::MemFunc) +// CHECK: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'void ()' +// CHECK-NEXT: NestedNameSpecifier{{.*}} 'S' +#pragma acc routine(S::StaticMemFunc) +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'void ()' +// CHECK-NEXT: NestedNameSpecifier{{.*}} 'S' +#pragma acc routine(S::Lambda) +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const S::(lambda at +// CHECK-NEXT: NestedNameSpecifier{{.*}} 'S' + +#pragma acc routine(MemFunc) +// CHECK: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'void ()' +#pragma acc routine(StaticMemFunc) +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'void ()' +#pragma acc routine(Lambda) +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const S::(lambda at +}; + +#pragma acc routine(S::MemFunc) +// CHECK: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'void ()' +// CHECK-NEXT: NestedNameSpecifier{{.*}} 'S' +#pragma acc routine(S::StaticMemFunc) +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'void ()' +// CHECK-NEXT: NestedNameSpecifier{{.*}} 'S' +#pragma acc routine(S::Lambda) +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const S::(lambda at +// CHECK-NEXT: NestedNameSpecifier{{.*}} 'S' + +template<typename T> +struct DepS { + T MemFunc(); + static T StaticMemFunc(); + constexpr static auto Lambda = [](){}; + +#pragma acc routine(Lambda) +// CHECK: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const auto' +#pragma acc routine(MemFunc) +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'T ()' +#pragma acc routine(StaticMemFunc) +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'T ()' + +#pragma acc routine(DepS::Lambda) +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const auto' +// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<T>' +#pragma acc routine(DepS::MemFunc) +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'T ()' +// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<T>' +#pragma acc routine(DepS::StaticMemFunc) +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'T ()' +// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<T>' + +#pragma acc routine(DepS<T>::Lambda) +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const auto' +// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<T>' +#pragma acc routine(DepS<T>::MemFunc) +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'T ()' +// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<T>' +#pragma acc routine(DepS<T>::StaticMemFunc) +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'T ()' +// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<T>' + +// Instantiation: +// CHECK: ClassTemplateSpecializationDecl{{.*}}struct DepS +// CHECK: CXXRecordDecl{{.*}} struct DepS + +// CHECK: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const DepS<int>:: + +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'int ()' + +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'int ()' + +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const DepS<int>:: +// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>' + +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'int ()' +// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>' + +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'int ()' +// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>' + +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const DepS<int>:: +// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>' + +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'int ()' +// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>' + +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'int ()' +// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>' +}; + +#pragma acc routine(DepS<int>::Lambda) +// CHECK: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const DepS<int>:: +// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>' +#pragma acc routine(DepS<int>::MemFunc) +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'int ()' +// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>' +#pragma acc routine(DepS<int>::StaticMemFunc) +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'int ()' +// CHECK-NEXT: NestedNameSpecifier{{.*}} 'DepS<int>' + +template<typename T> +void TemplFunc() { +#pragma acc routine(T::MemFunc) +// CHECK: DeclStmt +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DependentScopeDeclRefExpr{{.*}}'<dependent type>' +// CHECK-NEXT: NestedNameSpecifier{{.*}} 'T' +#pragma acc routine(T::StaticMemFunc) +// CHECK-NEXT: DeclStmt +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DependentScopeDeclRefExpr{{.*}}'<dependent type>' +// CHECK-NEXT: NestedNameSpecifier{{.*}} 'T' +#pragma acc routine(T::Lambda) +// CHECK-NEXT: DeclStmt +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DependentScopeDeclRefExpr{{.*}}'<dependent type>' +// CHECK-NEXT: NestedNameSpecifier{{.*}} 'T' + +// Instantiation: +// CHECK: FunctionDecl{{.*}} TemplFunc 'void ()' implicit_instantiation +// CHECK: DeclStmt +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'MemFunc' 'void ()' +// CHECK-NEXT: NestedNameSpecifier{{.*}} 'S' +// CHECK-NEXT: DeclStmt +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'StaticMemFunc' 'void ()' +// CHECK-NEXT: NestedNameSpecifier{{.*}} 'S' +// CHECK-NEXT: DeclStmt +// CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified +// CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' 'const S::(lambda at +// CHECK-NEXT: NestedNameSpecifier{{.*}} 'S' +} + +void usage() { + DepS<int> s; + TemplFunc<S>(); +} +#endif diff --git a/clang/test/SemaOpenACC/routine-construct.cpp b/clang/test/SemaOpenACC/routine-construct.cpp new file mode 100644 index 0000000..f953a45 --- /dev/null +++ b/clang/test/SemaOpenACC/routine-construct.cpp @@ -0,0 +1,400 @@ +// RUN: %clang_cc1 %s -fopenacc -verify + +// expected-error@+1{{use of undeclared identifier 'UnnamedYet'}} +#pragma acc routine(UnnamedYet) +void UnnamedYet(); +// expected-error@+1{{use of undeclared identifier 'Invalid'}} +#pragma acc routine(Invalid) + +// Fine, since these are the same function. +void SameFunc(); +#pragma acc routine(SameFunc) +void SameFunc(); + +void NoMagicStatic() { + static int F = 1; +} +// expected-error@-2{{function static variables are not permitted in functions to which an OpenACC 'routine' directive applies}} +// expected-note@+1{{'routine' construct is here}} +#pragma acc routine(NoMagicStatic) + +void NoMagicStatic2(); +// expected-error@+4{{function static variables are not permitted in functions to which an OpenACC 'routine' directive applies}} +// expected-note@+1{{'routine' construct is here}} +#pragma acc routine(NoMagicStatic2) +void NoMagicStatic2() { + static int F = 1; +} + +void HasMagicStaticLambda() { + auto MSLambda = []() { + static int I = 5; + }; +// expected-error@-2{{function static variables are not permitted in functions to which an OpenACC 'routine' directive applies}} +// expected-note@+1{{'routine' construct is here}} +#pragma acc routine (MSLambda) +} + +auto Lambda = [](){}; +#pragma acc routine(Lambda) +auto GenLambda = [](auto){}; +// expected-error@+1{{OpenACC routine name 'GenLambda' names a set of overloads}} +#pragma acc routine(GenLambda) +// Variable? +int Variable; +// Plain function +int function(); + +#pragma acc routine (function) +// expected-error@+1{{OpenACC routine name 'Variable' does not name a function}} +#pragma acc routine (Variable) + +// Var template? +template<typename T> +T VarTempl = 0; +// expected-error@+2{{use of variable template 'VarTempl' requires template arguments}} +// expected-note@-2{{template is declared here}} +#pragma acc routine (VarTempl) +// expected-error@+1{{OpenACC routine name 'VarTempl<int>' does not name a function}} +#pragma acc routine (VarTempl<int>) + +// Function in NS +namespace NS { + int NSFunc(); +auto Lambda = [](){}; +} +#pragma acc routine(NS::NSFunc) +#pragma acc routine(NS::Lambda) + +// Ambiguous Function +int ambig_func(); +int ambig_func(int); +// expected-error@+1{{OpenACC routine name 'ambig_func' names a set of overloads}} +#pragma acc routine (ambig_func) + +// Ambiguous in NS +namespace NS { +int ambig_func(); +int ambig_func(int); +} +// expected-error@+1{{OpenACC routine name 'NS::ambig_func' names a set of overloads}} +#pragma acc routine (NS::ambig_func) + +// function template +template<typename T, typename U> +void templ_func(); + +// expected-error@+1{{OpenACC routine name 'templ_func' names a set of overloads}} +#pragma acc routine(templ_func) +// expected-error@+1{{OpenACC routine name 'templ_func<int>' names a set of overloads}} +#pragma acc routine(templ_func<int>) +// expected-error@+1{{OpenACC routine name 'templ_func<int, float>' names a set of overloads}} +#pragma acc routine(templ_func<int, float>) + +struct S { + void MemFunc(); + static void StaticMemFunc(); + template<typename U> + void TemplMemFunc(); + template<typename U> + static void TemplStaticMemFunc(); + + void MemFuncAmbig(); + void MemFuncAmbig(int); + template<typename T> + void TemplMemFuncAmbig(); + template<typename T> + void TemplMemFuncAmbig(int); + + int Field; + + constexpr static auto Lambda = [](){}; + +#pragma acc routine(S::MemFunc) +#pragma acc routine(S::StaticMemFunc) +#pragma acc routine(S::Lambda) +// expected-error@+1{{OpenACC routine name 'S::TemplMemFunc' names a set of overloads}} +#pragma acc routine(S::TemplMemFunc) +// expected-error@+1{{OpenACC routine name 'S::TemplStaticMemFunc' names a set of overloads}} +#pragma acc routine(S::TemplStaticMemFunc) +// expected-error@+1{{OpenACC routine name 'S::template TemplMemFunc<int>' names a set of overloads}} +#pragma acc routine(S::template TemplMemFunc<int>) +// expected-error@+1{{OpenACC routine name 'S::template TemplStaticMemFunc<int>' names a set of overloads}} +#pragma acc routine(S::template TemplStaticMemFunc<int>) +// expected-error@+1{{OpenACC routine name 'S::MemFuncAmbig' names a set of overloads}} +#pragma acc routine(S::MemFuncAmbig) +// expected-error@+1{{OpenACC routine name 'S::TemplMemFuncAmbig' names a set of overloads}} +#pragma acc routine(S::TemplMemFuncAmbig) +// expected-error@+1{{OpenACC routine name 'S::template TemplMemFuncAmbig<int>' names a set of overloads}} +#pragma acc routine(S::template TemplMemFuncAmbig<int>) +// expected-error@+1{{OpenACC routine name 'S::Field' does not name a function}} +#pragma acc routine(S::Field) +}; + +#pragma acc routine(S::MemFunc) +#pragma acc routine(S::StaticMemFunc) +#pragma acc routine(S::Lambda) +// expected-error@+1{{OpenACC routine name 'S::TemplMemFunc' names a set of overloads}} +#pragma acc routine(S::TemplMemFunc) +// expected-error@+1{{OpenACC routine name 'S::TemplStaticMemFunc' names a set of overloads}} +#pragma acc routine(S::TemplStaticMemFunc) +// expected-error@+1{{OpenACC routine name 'S::template TemplMemFunc<int>' names a set of overloads}} +#pragma acc routine(S::template TemplMemFunc<int>) +// expected-error@+1{{OpenACC routine name 'S::template TemplStaticMemFunc<int>' names a set of overloads}} +#pragma acc routine(S::template TemplStaticMemFunc<int>) +// expected-error@+1{{OpenACC routine name 'S::MemFuncAmbig' names a set of overloads}} +#pragma acc routine(S::MemFuncAmbig) +// expected-error@+1{{OpenACC routine name 'S::TemplMemFuncAmbig' names a set of overloads}} +#pragma acc routine(S::TemplMemFuncAmbig) +// expected-error@+1{{OpenACC routine name 'S::template TemplMemFuncAmbig<int>' names a set of overloads}} +#pragma acc routine(S::template TemplMemFuncAmbig<int>) +// expected-error@+1{{OpenACC routine name 'S::Field' does not name a function}} +#pragma acc routine(S::Field) + +template<typename T> +struct DepS { // #DEPS + void MemFunc(); + static void StaticMemFunc(); + + template<typename U> + void TemplMemFunc(); + template<typename U> + static void TemplStaticMemFunc(); + + void MemFuncAmbig(); + void MemFuncAmbig(int); + template<typename U> + void TemplMemFuncAmbig(); + template<typename U> + void TemplMemFuncAmbig(int); + + int Field; + constexpr static auto Lambda = [](){}; + // expected-error@+2{{non-const static data member must be initialized out of line}} + // expected-note@#DEPSInst{{in instantiation of template class}} + static auto LambdaBroken = [](){}; + +#pragma acc routine(DepS::MemFunc) +#pragma acc routine(DepS::StaticMemFunc) +#pragma acc routine(DepS::Lambda) +#pragma acc routine(DepS::LambdaBroken) +// expected-error@+1{{OpenACC routine name 'DepS<T>::TemplMemFunc' names a set of overloads}} +#pragma acc routine(DepS::TemplMemFunc) +// expected-error@+1{{OpenACC routine name 'DepS<T>::TemplStaticMemFunc' names a set of overloads}} +#pragma acc routine(DepS::TemplStaticMemFunc) +// expected-error@+1{{OpenACC routine name 'DepS<T>::template TemplMemFunc<int>' names a set of overloads}} +#pragma acc routine(DepS::template TemplMemFunc<int>) +// expected-error@+1{{OpenACC routine name 'DepS<T>::template TemplStaticMemFunc<int>' names a set of overloads}} +#pragma acc routine(DepS::template TemplStaticMemFunc<int>) +// expected-error@+1{{OpenACC routine name 'DepS<T>::MemFuncAmbig' names a set of overloads}} +#pragma acc routine(DepS::MemFuncAmbig) +// expected-error@+1{{OpenACC routine name 'DepS<T>::TemplMemFuncAmbig' names a set of overloads}} +#pragma acc routine(DepS::TemplMemFuncAmbig) +// expected-error@+1{{OpenACC routine name 'DepS<T>::template TemplMemFuncAmbig<int>' names a set of overloads}} +#pragma acc routine(DepS::template TemplMemFuncAmbig<int>) +// expected-error@+1{{OpenACC routine name 'DepS<T>::Field' does not name a function}} +#pragma acc routine(DepS::Field) + +#pragma acc routine(DepS<T>::MemFunc) +#pragma acc routine(DepS<T>::StaticMemFunc) +#pragma acc routine(DepS<T>::Lambda) +#pragma acc routine(DepS<T>::LambdaBroken) +// expected-error@+1{{OpenACC routine name 'DepS<T>::TemplMemFunc' names a set of overloads}} +#pragma acc routine(DepS<T>::TemplMemFunc) +// expected-error@+1{{OpenACC routine name 'DepS<T>::TemplStaticMemFunc' names a set of overloads}} +#pragma acc routine(DepS<T>::TemplStaticMemFunc) +// expected-error@+1{{OpenACC routine name 'DepS<T>::template TemplMemFunc<int>' names a set of overloads}} +#pragma acc routine(DepS<T>::template TemplMemFunc<int>) +// expected-error@+1{{OpenACC routine name 'DepS<T>::template TemplStaticMemFunc<int>' names a set of overloads}} +#pragma acc routine(DepS<T>::template TemplStaticMemFunc<int>) +// expected-error@+1{{OpenACC routine name 'DepS<T>::MemFuncAmbig' names a set of overloads}} +#pragma acc routine(DepS<T>::MemFuncAmbig) +// expected-error@+1{{OpenACC routine name 'DepS<T>::TemplMemFuncAmbig' names a set of overloads}} +#pragma acc routine(DepS<T>::TemplMemFuncAmbig) +// expected-error@+1{{OpenACC routine name 'DepS<T>::template TemplMemFuncAmbig<int>' names a set of overloads}} +#pragma acc routine(DepS<T>::template TemplMemFuncAmbig<int>) +// expected-error@+1{{OpenACC routine name 'DepS<T>::Field' does not name a function}} +#pragma acc routine(DepS<T>::Field) +}; + +void Inst() { + DepS<int> S; // #DEPSInst +} + + +//expected-error@+2{{use of class template 'DepS' requires template arguments}} +// expected-note@#DEPS{{template is declared here}} +#pragma acc routine(DepS::Lambda) +//expected-error@+2{{use of class template 'DepS' requires template arguments}} +// expected-note@#DEPS{{template is declared here}} +#pragma acc routine(DepS::MemFunc) +//expected-error@+2{{use of class template 'DepS' requires template arguments}} +// expected-note@#DEPS{{template is declared here}} +#pragma acc routine(DepS::StaticMemFunc) +//expected-error@+2{{use of class template 'DepS' requires template arguments}} +// expected-note@#DEPS{{template is declared here}} +#pragma acc routine(DepS::TemplMemFunc) +//expected-error@+2{{use of class template 'DepS' requires template arguments}} +// expected-note@#DEPS{{template is declared here}} +#pragma acc routine(DepS::TemplStaticMemFunc) +//expected-error@+2{{use of class template 'DepS' requires template arguments}} +// expected-note@#DEPS{{template is declared here}} +#pragma acc routine(DepS::TemplMemFunc<int>) +//expected-error@+2{{use of class template 'DepS' requires template arguments}} +// expected-note@#DEPS{{template is declared here}} +#pragma acc routine(DepS::TemplStaticMemFunc<int>) +//expected-error@+2{{use of class template 'DepS' requires template arguments}} +// expected-note@#DEPS{{template is declared here}} +#pragma acc routine(DepS::MemFuncAmbig) +//expected-error@+2{{use of class template 'DepS' requires template arguments}} +// expected-note@#DEPS{{template is declared here}} +#pragma acc routine(DepS::TemplMemFuncAmbig) +//expected-error@+2{{use of class template 'DepS' requires template arguments}} +// expected-note@#DEPS{{template is declared here}} +#pragma acc routine(DepS::TemplMemFuncAmbig<int>) +//expected-error@+2{{use of class template 'DepS' requires template arguments}} +// expected-note@#DEPS{{template is declared here}} +#pragma acc routine(DepS::Field) + +//expected-error@+1{{use of undeclared identifier 'T'}} +#pragma acc routine(DepS<T>::Lambda) +//expected-error@+1{{use of undeclared identifier 'T'}} +#pragma acc routine(DepS<T>::MemFunc) +//expected-error@+1{{use of undeclared identifier 'T'}} +#pragma acc routine(DepS<T>::StaticMemFunc) +//expected-error@+1{{use of undeclared identifier 'T'}} +#pragma acc routine(DepS<T>::TemplMemFunc) +//expected-error@+1{{use of undeclared identifier 'T'}} +#pragma acc routine(DepS<T>::TemplStaticMemFunc) +//expected-error@+1{{use of undeclared identifier 'T'}} +#pragma acc routine(DepS<T>::TemplMemFunc<int>) +//expected-error@+1{{use of undeclared identifier 'T'}} +#pragma acc routine(DepS<T>::TemplStaticMemFunc<int>) +//expected-error@+1{{use of undeclared identifier 'T'}} +#pragma acc routine(DepS<T>::MemFuncAmbig) +//expected-error@+1{{use of undeclared identifier 'T'}} +#pragma acc routine(DepS<T>::TemplMemFuncAmbig) +//expected-error@+1{{use of undeclared identifier 'T'}} +#pragma acc routine(DepS<T>::TemplMemFuncAmbig<int>) +//expected-error@+1{{use of undeclared identifier 'T'}} +#pragma acc routine(DepS<T>::Field) + +#pragma acc routine(DepS<int>::Lambda) +#pragma acc routine(DepS<int>::LambdaBroken) +#pragma acc routine(DepS<int>::MemFunc) +#pragma acc routine(DepS<int>::StaticMemFunc) +// expected-error@+1{{OpenACC routine name 'DepS<int>::TemplMemFunc' names a set of overloads}} +#pragma acc routine(DepS<int>::TemplMemFunc) +// expected-error@+1{{OpenACC routine name 'DepS<int>::TemplStaticMemFunc' names a set of overloads}} +#pragma acc routine(DepS<int>::TemplStaticMemFunc) +// expected-error@+1{{OpenACC routine name 'DepS<int>::TemplMemFunc<int>' names a set of overloads}} +#pragma acc routine(DepS<int>::TemplMemFunc<int>) +// expected-error@+1{{OpenACC routine name 'DepS<int>::TemplStaticMemFunc<int>' names a set of overloads}} +#pragma acc routine(DepS<int>::TemplStaticMemFunc<int>) +// expected-error@+1{{OpenACC routine name 'DepS<int>::MemFuncAmbig' names a set of overloads}} +#pragma acc routine(DepS<int>::MemFuncAmbig) +// expected-error@+1{{OpenACC routine name 'DepS<int>::TemplMemFuncAmbig' names a set of overloads}} +#pragma acc routine(DepS<int>::TemplMemFuncAmbig) +// expected-error@+1{{OpenACC routine name 'DepS<int>::TemplMemFuncAmbig<int>' names a set of overloads}} +#pragma acc routine(DepS<int>::TemplMemFuncAmbig<int>) +// expected-error@+1{{OpenACC routine name 'DepS<int>::Field' does not name a function}} +#pragma acc routine(DepS<int>::Field) + +template<typename T> +void TemplFunc() { +#pragma acc routine(T::MemFunc) +#pragma acc routine(T::StaticMemFunc) +#pragma acc routine(T::Lambda) +// expected-error@+1{{OpenACC routine name 'S::TemplMemFunc' names a set of overloads}} +#pragma acc routine(T::TemplMemFunc) +// expected-error@+1{{OpenACC routine name 'S::TemplStaticMemFunc' names a set of overloads}} +#pragma acc routine(T::TemplStaticMemFunc) +// expected-error@+1{{OpenACC routine name 'S::template TemplMemFunc<int>' names a set of overloads}} +#pragma acc routine(T::template TemplMemFunc<int>) +// expected-error@+1{{OpenACC routine name 'S::template TemplStaticMemFunc<int>' names a set of overloads}} +#pragma acc routine(T::template TemplStaticMemFunc<int>) +// expected-error@+1{{OpenACC routine name 'S::MemFuncAmbig' names a set of overloads}} +#pragma acc routine(T::MemFuncAmbig) +// expected-error@+1{{OpenACC routine name 'S::TemplMemFuncAmbig' names a set of overloads}} +#pragma acc routine(T::TemplMemFuncAmbig) +// expected-error@+1{{OpenACC routine name 'S::template TemplMemFuncAmbig<int>' names a set of overloads}} +#pragma acc routine(T::template TemplMemFuncAmbig<int>) +// expected-error@+1{{OpenACC routine name 'S::Field' does not name a function}} +#pragma acc routine(T::Field) +} + +template <typename T> +struct DepRefersToT { +#pragma acc routine(T::MemFunc) +#pragma acc routine(T::StaticMemFunc) +#pragma acc routine(T::Lambda) +// expected-error@+1{{OpenACC routine name 'S::TemplMemFunc' names a set of overloads}} +#pragma acc routine(T::TemplMemFunc) +// expected-error@+1{{OpenACC routine name 'S::TemplStaticMemFunc' names a set of overloads}} +#pragma acc routine(T::TemplStaticMemFunc) +// expected-error@+1{{OpenACC routine name 'S::template TemplMemFunc<int>' names a set of overloads}} +#pragma acc routine(T::template TemplMemFunc<int>) +// expected-error@+1{{OpenACC routine name 'S::template TemplStaticMemFunc<int>' names a set of overloads}} +#pragma acc routine(T::template TemplStaticMemFunc<int>) +// expected-error@+1{{OpenACC routine name 'S::MemFuncAmbig' names a set of overloads}} +#pragma acc routine(T::MemFuncAmbig) +// expected-error@+1{{OpenACC routine name 'S::TemplMemFuncAmbig' names a set of overloads}} +#pragma acc routine(T::TemplMemFuncAmbig) +// expected-error@+1{{OpenACC routine name 'S::template TemplMemFuncAmbig<int>' names a set of overloads}} +#pragma acc routine(T::template TemplMemFuncAmbig<int>) +// expected-error@+1{{OpenACC routine name 'S::Field' does not name a function}} +#pragma acc routine(T::Field) + + void MemFunc() { +#pragma acc routine(T::MemFunc) +#pragma acc routine(T::StaticMemFunc) +#pragma acc routine(T::Lambda) +// expected-error@+1{{OpenACC routine name 'S::TemplMemFunc' names a set of overloads}} +#pragma acc routine(T::TemplMemFunc) +// expected-error@+1{{OpenACC routine name 'S::TemplStaticMemFunc' names a set of overloads}} +#pragma acc routine(T::TemplStaticMemFunc) +// expected-error@+1{{OpenACC routine name 'S::template TemplMemFunc<int>' names a set of overloads}} +#pragma acc routine(T::template TemplMemFunc<int>) +// expected-error@+1{{OpenACC routine name 'S::template TemplStaticMemFunc<int>' names a set of overloads}} +#pragma acc routine(T::template TemplStaticMemFunc<int>) +// expected-error@+1{{OpenACC routine name 'S::MemFuncAmbig' names a set of overloads}} +#pragma acc routine(T::MemFuncAmbig) +// expected-error@+1{{OpenACC routine name 'S::TemplMemFuncAmbig' names a set of overloads}} +#pragma acc routine(T::TemplMemFuncAmbig) +// expected-error@+1{{OpenACC routine name 'S::template TemplMemFuncAmbig<int>' names a set of overloads}} +#pragma acc routine(T::template TemplMemFuncAmbig<int>) +// expected-error@+1{{OpenACC routine name 'S::Field' does not name a function}} +#pragma acc routine(T::Field) + } + + template<typename U> + void TemplMemFunc() { +#pragma acc routine(T::MemFunc) +#pragma acc routine(T::StaticMemFunc) +#pragma acc routine(T::Lambda) +// expected-error@+1{{OpenACC routine name 'S::TemplMemFunc' names a set of overloads}} +#pragma acc routine(T::TemplMemFunc) +// expected-error@+1{{OpenACC routine name 'S::TemplStaticMemFunc' names a set of overloads}} +#pragma acc routine(T::TemplStaticMemFunc) +// expected-error@+1{{OpenACC routine name 'S::template TemplMemFunc<int>' names a set of overloads}} +#pragma acc routine(T::template TemplMemFunc<int>) +// expected-error@+1{{OpenACC routine name 'S::template TemplStaticMemFunc<int>' names a set of overloads}} +#pragma acc routine(T::template TemplStaticMemFunc<int>) +// expected-error@+1{{OpenACC routine name 'S::MemFuncAmbig' names a set of overloads}} +#pragma acc routine(T::MemFuncAmbig) +// expected-error@+1{{OpenACC routine name 'S::TemplMemFuncAmbig' names a set of overloads}} +#pragma acc routine(T::TemplMemFuncAmbig) +// expected-error@+1{{OpenACC routine name 'S::template TemplMemFuncAmbig<int>' names a set of overloads}} +#pragma acc routine(T::template TemplMemFuncAmbig<int>) +// expected-error@+1{{OpenACC routine name 'S::Field' does not name a function}} +#pragma acc routine(T::Field) + } + +}; + +void inst() { + TemplFunc<S>(); // expected-note{{in instantiation of}} + DepRefersToT<S> s; // expected-note{{in instantiation of}} + s.MemFunc(); // expected-note{{in instantiation of}} + s.TemplMemFunc<S>(); // expected-note{{in instantiation of}} +} diff --git a/clang/test/SemaOpenACC/unimplemented-construct.c b/clang/test/SemaOpenACC/unimplemented-construct.c index 261c855..ec11792 100644 --- a/clang/test/SemaOpenACC/unimplemented-construct.c +++ b/clang/test/SemaOpenACC/unimplemented-construct.c @@ -1,39 +1,38 @@ // RUN: %clang_cc1 %s -verify -fopenacc -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} +// expected-warning@+1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}} #pragma acc routine struct S { -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} +// expected-warning@+1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}} #pragma acc routine int foo; }; void func() { -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} +// expected-warning@+1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}} #pragma acc routine int foo; -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} +// expected-warning@+1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}} #pragma acc routine { -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} +// expected-warning@+1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}} #pragma acc routine { -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} +// expected-warning@+1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}} #pragma acc routine } } -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} +// expected-warning@+1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}} #pragma acc routine while(0){} -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} +// expected-warning@+1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}} #pragma acc routine for(;;){} -// expected-warning@+1{{OpenACC construct 'routine' not yet implemented, pragma ignored}} +// expected-warning@+1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}} #pragma acc routine }; - diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp index dba5516..e536fa6aa 100644 --- a/clang/tools/libclang/CIndex.cpp +++ b/clang/tools/libclang/CIndex.cpp @@ -7265,6 +7265,7 @@ CXCursor clang_getCursorDefinition(CXCursor C) { case Decl::RequiresExprBody: case Decl::UnresolvedUsingIfExists: case Decl::OpenACCDeclare: + case Decl::OpenACCRoutine: return C; // Declaration kinds that don't make any sense here, but are |