diff options
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 |