Author: erichkeane Date: 2025-03-10T07:49:13-07:00 New Revision: 8a8f1359ee1d47d85c5fb4d23587845baecd59c9
URL: https://github.com/llvm/llvm-project/commit/8a8f1359ee1d47d85c5fb4d23587845baecd59c9 DIFF: https://github.com/llvm/llvm-project/commit/8a8f1359ee1d47d85c5fb4d23587845baecd59c9.diff LOG: [OpenACC] Implement 'bind' ast/sema for 'routine' directive The 'bind' clause allows the renaming of a function during code generation. There are a few rules about when this can/cannot happen, and it takes either a string or identifier (previously mis-implemetned as ID-expression) argument. Note there are additional rules to this in the implicit-function routine case, but that isn't implemented in this patch, as implicit-function routine is not yet implemented either. Added: Modified: clang/include/clang/AST/OpenACCClause.h clang/include/clang/Basic/Attr.td clang/include/clang/Basic/DiagnosticSemaKinds.td clang/include/clang/Basic/OpenACCClauses.def clang/include/clang/Parse/Parser.h clang/include/clang/Sema/SemaOpenACC.h clang/lib/AST/OpenACCClause.cpp clang/lib/AST/StmtProfile.cpp clang/lib/AST/TextNodeDumper.cpp clang/lib/Parse/ParseOpenACC.cpp clang/lib/Sema/SemaOpenACC.cpp clang/lib/Sema/SemaOpenACCClause.cpp clang/lib/Sema/SemaTemplateInstantiateDecl.cpp clang/lib/Sema/TreeTransform.h clang/lib/Serialization/ASTReader.cpp clang/lib/Serialization/ASTWriter.cpp clang/test/AST/ast-print-openacc-routine-construct.cpp clang/test/ParserOpenACC/parse-clauses.c clang/test/ParserOpenACC/parse-clauses.cpp clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c clang/test/SemaOpenACC/combined-construct-device_type-clause.c clang/test/SemaOpenACC/compute-construct-device_type-clause.c clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c clang/test/SemaOpenACC/loop-construct-device_type-clause.c clang/test/SemaOpenACC/routine-construct-ast.cpp clang/test/SemaOpenACC/routine-construct-clauses.cpp clang/tools/libclang/CIndex.cpp Removed: ################################################################################ diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h index b2cf621bc0a78..049389229b8d5 100644 --- a/clang/include/clang/AST/OpenACCClause.h +++ b/clang/include/clang/AST/OpenACCClause.h @@ -18,6 +18,7 @@ #include "clang/Basic/OpenACCKinds.h" #include <utility> +#include <variant> namespace clang { /// This is the base type for all OpenACC Clauses. @@ -206,6 +207,50 @@ class OpenACCClauseWithParams : public OpenACCClause { } }; +class OpenACCBindClause final : public OpenACCClauseWithParams { + std::variant<const StringLiteral *, const IdentifierInfo *> Argument; + + OpenACCBindClause(SourceLocation BeginLoc, SourceLocation LParenLoc, + const clang::StringLiteral *SL, SourceLocation EndLoc) + : OpenACCClauseWithParams(OpenACCClauseKind::Bind, BeginLoc, LParenLoc, + EndLoc), + Argument(SL) {} + OpenACCBindClause(SourceLocation BeginLoc, SourceLocation LParenLoc, + const IdentifierInfo *ID, SourceLocation EndLoc) + : OpenACCClauseWithParams(OpenACCClauseKind::Bind, BeginLoc, LParenLoc, + EndLoc), + Argument(ID) {} + +public: + static bool classof(const OpenACCClause *C) { + return C->getClauseKind() == OpenACCClauseKind::Bind; + } + static OpenACCBindClause *Create(const ASTContext &C, SourceLocation BeginLoc, + SourceLocation LParenLoc, + const IdentifierInfo *ID, + SourceLocation EndLoc); + static OpenACCBindClause *Create(const ASTContext &C, SourceLocation BeginLoc, + SourceLocation LParenLoc, + const StringLiteral *SL, + SourceLocation EndLoc); + + bool isStringArgument() const { + return std::holds_alternative<const StringLiteral *>(Argument); + } + + const StringLiteral *getStringArgument() const { + return std::get<const StringLiteral *>(Argument); + } + + bool isIdentifierArgument() const { + return std::holds_alternative<const IdentifierInfo *>(Argument); + } + + const IdentifierInfo *getIdentifierArgument() const { + return std::get<const IdentifierInfo *>(Argument); + } +}; + using DeviceTypeArgument = std::pair<IdentifierInfo *, SourceLocation>; /// A 'device_type' or 'dtype' clause, takes a list of either an 'asterisk' or /// an identifier. The 'asterisk' means 'the rest'. diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 8d7a9ef0cb1cd..4b4337cf460f3 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -5026,4 +5026,7 @@ def OpenACCRoutineAnnot : InheritableAttr { let Spellings = []; let Subjects = SubjectList<[Function]>; let Documentation = [InternalOnly]; + let AdditionalMembers = [{ + SourceLocation BindClause; + }]; } diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index eeb7e236e8b7a..b73ac2933b1ca 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12844,9 +12844,6 @@ 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">, - InGroup<SourceUsesOpenACC>; def err_acc_construct_appertainment : Error<"OpenACC construct '%0' cannot be used here; it can only " "be used in a statement context">; @@ -13087,6 +13084,9 @@ def err_acc_routine_overload_set def err_acc_magic_static_in_routine : Error<"function static variables are not permitted in functions to which " "an OpenACC 'routine' directive applies">; +def err_acc_duplicate_bind + : Error<"multiple 'routine' directives with 'bind' clauses are not " + "permitted to refer to the same function">; // AMDGCN builtins diagnostics def err_amdgcn_global_load_lds_size_invalid_value : Error<"invalid size value">; diff --git a/clang/include/clang/Basic/OpenACCClauses.def b/clang/include/clang/Basic/OpenACCClauses.def index f04965363f25e..a19ee69d3d190 100644 --- a/clang/include/clang/Basic/OpenACCClauses.def +++ b/clang/include/clang/Basic/OpenACCClauses.def @@ -24,6 +24,7 @@ VISIT_CLAUSE(Auto) VISIT_CLAUSE(Async) VISIT_CLAUSE(Attach) +VISIT_CLAUSE(Bind) VISIT_CLAUSE(Collapse) VISIT_CLAUSE(Copy) CLAUSE_ALIAS(PCopy, Copy, true) diff --git a/clang/include/clang/Parse/Parser.h b/clang/include/clang/Parse/Parser.h index 049156e266c70..1c8caf55a546d 100644 --- a/clang/include/clang/Parse/Parser.h +++ b/clang/include/clang/Parse/Parser.h @@ -3783,8 +3783,9 @@ class Parser : public CodeCompletionHandler { OpenACCWaitParseInfo ParseOpenACCWaitArgument(SourceLocation Loc, bool IsDirective); /// Parses the clause of the 'bind' argument, which can be a string literal or - /// an ID expression. - ExprResult ParseOpenACCBindClauseArgument(); + /// an identifier. + std::variant<std::monostate, StringLiteral *, IdentifierInfo *> + ParseOpenACCBindClauseArgument(); /// A type to represent the state of parsing after an attempt to parse an /// OpenACC int-expr. This is useful to determine whether an int-expr list can diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h index 748dcdd251a92..358dec4cadb72 100644 --- a/clang/include/clang/Sema/SemaOpenACC.h +++ b/clang/include/clang/Sema/SemaOpenACC.h @@ -261,10 +261,14 @@ class SemaOpenACC : public SemaBase { SmallVector<OpenACCGangKind> GangKinds; SmallVector<Expr *> IntExprs; }; + struct BindDetails { + std::variant<std::monostate, clang::StringLiteral *, IdentifierInfo *> + Argument; + }; std::variant<std::monostate, DefaultDetails, ConditionDetails, IntExprDetails, VarListDetails, WaitDetails, DeviceTypeDetails, - ReductionDetails, CollapseDetails, GangDetails> + ReductionDetails, CollapseDetails, GangDetails, BindDetails> Details = std::monostate{}; public: @@ -468,6 +472,13 @@ class SemaOpenACC : public SemaBase { return std::get<DeviceTypeDetails>(Details).Archs; } + std::variant<std::monostate, clang::StringLiteral *, IdentifierInfo *> + getBindDetails() const { + assert(ClauseKind == OpenACCClauseKind::Bind && + "Only 'bind' has bind details"); + return std::get<BindDetails>(Details).Argument; + } + void setLParenLoc(SourceLocation EndLoc) { LParenLoc = EndLoc; } void setEndLoc(SourceLocation EndLoc) { ClauseRange.setEnd(EndLoc); } @@ -652,6 +663,14 @@ class SemaOpenACC : public SemaBase { "Only 'collapse' has collapse details"); Details = CollapseDetails{IsForce, LoopCount}; } + + void setBindDetails( + std::variant<std::monostate, clang::StringLiteral *, IdentifierInfo *> + Arg) { + assert(ClauseKind == OpenACCClauseKind::Bind && + "Only 'bind' has bind details"); + Details = BindDetails{Arg}; + } }; SemaOpenACC(Sema &S); diff --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp index fd2c38a0e64e7..579ee91b7ec99 100644 --- a/clang/lib/AST/OpenACCClause.cpp +++ b/clang/lib/AST/OpenACCClause.cpp @@ -20,7 +20,8 @@ using namespace clang; bool OpenACCClauseWithParams::classof(const OpenACCClause *C) { return OpenACCDeviceTypeClause::classof(C) || OpenACCClauseWithCondition::classof(C) || - OpenACCClauseWithExprs::classof(C) || OpenACCSelfClause::classof(C); + OpenACCBindClause::classof(C) || OpenACCClauseWithExprs::classof(C) || + OpenACCSelfClause::classof(C); } bool OpenACCClauseWithExprs::classof(const OpenACCClause *C) { return OpenACCWaitClause::classof(C) || OpenACCNumGangsClause::classof(C) || @@ -609,6 +610,24 @@ OpenACCIfPresentClause *OpenACCIfPresentClause::Create(const ASTContext &C, return new (Mem) OpenACCIfPresentClause(BeginLoc, EndLoc); } +OpenACCBindClause *OpenACCBindClause::Create(const ASTContext &C, + SourceLocation BeginLoc, + SourceLocation LParenLoc, + const StringLiteral *SL, + SourceLocation EndLoc) { + void *Mem = C.Allocate(sizeof(OpenACCBindClause), alignof(OpenACCBindClause)); + return new (Mem) OpenACCBindClause(BeginLoc, LParenLoc, SL, EndLoc); +} + +OpenACCBindClause *OpenACCBindClause::Create(const ASTContext &C, + SourceLocation BeginLoc, + SourceLocation LParenLoc, + const IdentifierInfo *ID, + SourceLocation EndLoc) { + void *Mem = C.Allocate(sizeof(OpenACCBindClause), alignof(OpenACCBindClause)); + return new (Mem) OpenACCBindClause(BeginLoc, LParenLoc, ID, EndLoc); +} + //===----------------------------------------------------------------------===// // OpenACC clauses printing methods //===----------------------------------------------------------------------===// @@ -936,3 +955,12 @@ void OpenACCClausePrinter::VisitIfPresentClause( const OpenACCIfPresentClause &C) { OS << "if_present"; } + +void OpenACCClausePrinter::VisitBindClause(const OpenACCBindClause &C) { + OS << "bind("; + if (C.isStringArgument()) + OS << '"' << C.getStringArgument()->getString() << '"'; + else + OS << C.getIdentifierArgument()->getName(); + OS << ")"; +} diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp index 574f67f4274e7..b543073de24bb 100644 --- a/clang/lib/AST/StmtProfile.cpp +++ b/clang/lib/AST/StmtProfile.cpp @@ -2732,6 +2732,10 @@ void OpenACCClauseProfiler::VisitReductionClause( const OpenACCReductionClause &Clause) { VisitClauseWithVarList(Clause); } + +void OpenACCClauseProfiler::VisitBindClause(const OpenACCBindClause &Clause) { + assert(false && "not implemented... what can we do about our expr?"); +} } // namespace void StmtProfiler::VisitOpenACCComputeConstruct( diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp index 91f3f14c6b454..7600b40a7482e 100644 --- a/clang/lib/AST/TextNodeDumper.cpp +++ b/clang/lib/AST/TextNodeDumper.cpp @@ -435,6 +435,7 @@ void TextNodeDumper::Visit(const OpenACCClause *C) { case OpenACCClauseKind::UseDevice: case OpenACCClauseKind::Vector: case OpenACCClauseKind::VectorLength: + case OpenACCClauseKind::Invalid: // The condition expression will be printed as a part of the 'children', // but print 'clause' here so it is clear what is happening from the dump. OS << " clause"; @@ -501,9 +502,15 @@ void TextNodeDumper::Visit(const OpenACCClause *C) { OS << " clause Operator: " << cast<OpenACCReductionClause>(C)->getReductionOp(); break; - default: - // Nothing to do here. - break; + case OpenACCClauseKind::Bind: + OS << " clause"; + if (cast<OpenACCBindClause>(C)->isIdentifierArgument()) + OS << " identifier '" + << cast<OpenACCBindClause>(C)->getIdentifierArgument()->getName() + << "'"; + else + AddChild( + [=] { Visit(cast<OpenACCBindClause>(C)->getStringArgument()); }); } } dumpPointer(C); diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp index 6ea17c97d6345..8d3501082cc27 100644 --- a/clang/lib/Parse/ParseOpenACC.cpp +++ b/clang/lib/Parse/ParseOpenACC.cpp @@ -1059,8 +1059,11 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams( break; } case OpenACCClauseKind::Bind: { - ExprResult BindArg = ParseOpenACCBindClauseArgument(); - if (BindArg.isInvalid()) { + ParsedClause.setBindDetails(ParseOpenACCBindClauseArgument()); + + // We can create an 'empty' bind clause in the event of an error + if (std::holds_alternative<std::monostate>( + ParsedClause.getBindDetails())) { Parens.skipToEnd(); return OpenACCCanContinue(); } @@ -1334,7 +1337,8 @@ ExprResult Parser::ParseOpenACCIDExpression() { return getActions().CorrectDelayedTyposInExpr(Res); } -ExprResult Parser::ParseOpenACCBindClauseArgument() { +std::variant<std::monostate, clang::StringLiteral *, IdentifierInfo *> +Parser::ParseOpenACCBindClauseArgument() { // OpenACC 3.3 section 2.15: // The bind clause specifies the name to use when calling the procedure on a // device other than the host. If the name is specified as an identifier, it @@ -1343,14 +1347,21 @@ ExprResult Parser::ParseOpenACCBindClauseArgument() { // name unmodified. if (getCurToken().is(tok::r_paren)) { Diag(getCurToken(), diag::err_acc_incorrect_bind_arg); - return ExprError(); + return std::monostate{}; } - if (tok::isStringLiteral(getCurToken().getKind())) - return getActions().CorrectDelayedTyposInExpr(ParseStringLiteralExpression( - /*AllowUserDefinedLiteral=*/false, /*Unevaluated=*/true)); + if (getCurToken().is(tok::identifier)) { + IdentifierInfo *II = getCurToken().getIdentifierInfo(); + ConsumeToken(); + return II; + } - return ParseOpenACCIDExpression(); + ExprResult Res = + getActions().CorrectDelayedTyposInExpr(ParseStringLiteralExpression( + /*AllowUserDefinedLiteral=*/false, /*Unevaluated=*/true)); + if (!Res.isUsable()) + return std::monostate{}; + return cast<StringLiteral>(Res.get()); } /// OpenACC 3.3, section 1.6: diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp index 4fe6bf5099a64..ec9e9527dca6f 100644 --- a/clang/lib/Sema/SemaOpenACC.cpp +++ b/clang/lib/Sema/SemaOpenACC.cpp @@ -1904,7 +1904,31 @@ DeclGroupRef SemaOpenACC::ActOnEndDeclDirective( Diag(DirLoc, diag::note_acc_construct_here) << OpenACCDirectiveKind::Routine; } - FD->addAttr(OpenACCRoutineAnnotAttr::Create(getASTContext(), DirLoc)); + + // OpenACC 3.3 2.15: + // A bind clause may not bind to a routine name that has a visible bind + // clause. + // TODO OpenACC: There is an exception to this rule that if these are the + // implicit function style (that is, without a name), they may have + // duplicates as long as they have the same name. + auto BindItr = llvm::find_if(Clauses, llvm::IsaPred<OpenACCBindClause>); + if (auto *A = FD->getAttr<OpenACCRoutineAnnotAttr>()) { + if (BindItr != Clauses.end()) { + if (A->BindClause.isInvalid()) { + // If we have a bind clause, and the function doesn't have one + // annotated yet, set it. + A->BindClause = (*BindItr)->getBeginLoc(); + } else { + Diag((*BindItr)->getBeginLoc(), diag::err_acc_duplicate_bind); + Diag(A->BindClause, diag::note_acc_previous_clause_here); + } + } + } else { + auto *RAA = OpenACCRoutineAnnotAttr::Create(getASTContext(), DirLoc); + FD->addAttr(RAA); + if (BindItr != Clauses.end()) + RAA->BindClause = (*BindItr)->getBeginLoc(); + } } return DeclGroupRef{RoutineDecl}; diff --git a/clang/lib/Sema/SemaOpenACCClause.cpp b/clang/lib/Sema/SemaOpenACCClause.cpp index 1d70178d03611..ad54e2bbe9495 100644 --- a/clang/lib/Sema/SemaOpenACCClause.cpp +++ b/clang/lib/Sema/SemaOpenACCClause.cpp @@ -483,6 +483,14 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind, return false; } } + case OpenACCClauseKind::Bind: { + switch (DirectiveKind) { + case OpenACCDirectiveKind::Routine: + return true; + default: + return false; + } + } } default: @@ -677,10 +685,6 @@ class SemaOpenACCClauseVisitor { SemaOpenACCClauseVisitor(SemaOpenACC &S, ArrayRef<const OpenACCClause *> ExistingClauses) : SemaRef(S), Ctx(S.getASTContext()), ExistingClauses(ExistingClauses) {} - // Once we've implemented everything, we shouldn't need this infrastructure. - // But in the meantime, we use this to help decide whether the clause was - // handled for this directive. - bool diagNotImplemented() { return NotImplemented; } OpenACCClause *Visit(SemaOpenACC::OpenACCParsedClause &Clause) { switch (Clause.getClauseKind()) { @@ -1985,6 +1989,17 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitCollapseClause( LoopCount.get(), Clause.getEndLoc()); } +OpenACCClause *SemaOpenACCClauseVisitor::VisitBindClause( + SemaOpenACC::OpenACCParsedClause &Clause) { + if (std::holds_alternative<StringLiteral *>(Clause.getBindDetails())) + return OpenACCBindClause::Create( + Ctx, Clause.getBeginLoc(), Clause.getLParenLoc(), + std::get<StringLiteral *>(Clause.getBindDetails()), Clause.getEndLoc()); + return OpenACCBindClause::Create( + Ctx, Clause.getBeginLoc(), Clause.getLParenLoc(), + std::get<IdentifierInfo *>(Clause.getBindDetails()), Clause.getEndLoc()); +} + // Return true if the two vars refer to the same variable, for the purposes of // equality checking. bool areVarsEqual(Expr *VarExpr1, Expr *VarExpr2) { @@ -2073,12 +2088,7 @@ SemaOpenACC::ActOnClause(ArrayRef<const OpenACCClause *> ExistingClauses, assert((!Result || Result->getClauseKind() == Clause.getClauseKind()) && "Created wrong clause?"); - if (Visitor.diagNotImplemented()) - Diag(Clause.getBeginLoc(), diag::warn_acc_clause_unimplemented) - << Clause.getClauseKind(); - return Result; - } /// OpenACC 3.3 section 2.5.15: diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp index a231ad72a9e7c..170c0b0d39f86 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -1249,6 +1249,21 @@ void OpenACCDeclClauseInstantiator::VisitDevicePtrClause( ParsedClause.getEndLoc()); } +void OpenACCDeclClauseInstantiator::VisitBindClause( + const OpenACCBindClause &C) { + // Nothing to instantiate, we support only string literal or identifier. + if (C.isStringArgument()) + NewClause = OpenACCBindClause::Create( + SemaRef.getASTContext(), ParsedClause.getBeginLoc(), + ParsedClause.getLParenLoc(), C.getStringArgument(), + ParsedClause.getEndLoc()); + else + NewClause = OpenACCBindClause::Create( + SemaRef.getASTContext(), ParsedClause.getBeginLoc(), + ParsedClause.getLParenLoc(), C.getIdentifierArgument(), + ParsedClause.getEndLoc()); +} + llvm::SmallVector<OpenACCClause *> InstantiateOpenACCClauseList( Sema &S, const MultiLevelTemplateArgumentList &MLTAL, OpenACCDirectiveKind DK, ArrayRef<const OpenACCClause *> ClauseList) { diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h index 9591fd4cfcc1c..f532cea245e54 100644 --- a/clang/lib/Sema/TreeTransform.h +++ b/clang/lib/Sema/TreeTransform.h @@ -11878,7 +11878,12 @@ void OpenACCClauseTransform<Derived>::VisitDeviceResidentClause( template <typename Derived> void OpenACCClauseTransform<Derived>::VisitNoHostClause( const OpenACCNoHostClause &C) { - llvm_unreachable("device_resident clause not valid unless a decl transform"); + llvm_unreachable("nohost clause not valid unless a decl transform"); +} +template <typename Derived> +void OpenACCClauseTransform<Derived>::VisitBindClause( + const OpenACCBindClause &C) { + llvm_unreachable("bind clause not valid unless a decl transform"); } template <typename Derived> diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp index 2ac9754f02eed..590c7b4e40bff 100644 --- a/clang/lib/Serialization/ASTReader.cpp +++ b/clang/lib/Serialization/ASTReader.cpp @@ -12797,7 +12797,15 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() { LParenLoc, VarList, EndLoc); } - case OpenACCClauseKind::Bind: + case OpenACCClauseKind::Bind: { + SourceLocation LParenLoc = readSourceLocation(); + bool IsString = readBool(); + if (IsString) + return OpenACCBindClause::Create(getContext(), BeginLoc, LParenLoc, + cast<StringLiteral>(readExpr()), EndLoc); + return OpenACCBindClause::Create(getContext(), BeginLoc, LParenLoc, + readIdentifier(), EndLoc); + } case OpenACCClauseKind::Invalid: llvm_unreachable("Clause serialization not yet implemented"); } diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp index 0aa115ecadf8e..d7328ccc8f6c6 100644 --- a/clang/lib/Serialization/ASTWriter.cpp +++ b/clang/lib/Serialization/ASTWriter.cpp @@ -8844,7 +8844,17 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) { return; } - case OpenACCClauseKind::Bind: + case OpenACCClauseKind::Bind: { + const auto *BC = cast<OpenACCBindClause>(C); + writeSourceLocation(BC->getLParenLoc()); + writeBool(BC->isStringArgument()); + if (BC->isStringArgument()) + AddStmt(const_cast<StringLiteral *>(BC->getStringArgument())); + else + AddIdentifierRef(BC->getIdentifierArgument()); + + return; + } case OpenACCClauseKind::Invalid: llvm_unreachable("Clause serialization not yet implemented"); } diff --git a/clang/test/AST/ast-print-openacc-routine-construct.cpp b/clang/test/AST/ast-print-openacc-routine-construct.cpp index df41c61142d37..1d21e2ddee9a1 100644 --- a/clang/test/AST/ast-print-openacc-routine-construct.cpp +++ b/clang/test/AST/ast-print-openacc-routine-construct.cpp @@ -1,11 +1,11 @@ // RUN: %clang_cc1 -fopenacc -ast-print %s -o - | FileCheck %s auto Lambda = [](){}; -// CHECK: #pragma acc routine(Lambda) worker -#pragma acc routine(Lambda) worker +// CHECK: #pragma acc routine(Lambda) worker bind(identifier) +#pragma acc routine(Lambda) worker bind(identifier) int function(); -// CHECK: #pragma acc routine(function) vector nohost -#pragma acc routine (function) vector nohost +// CHECK: #pragma acc routine(function) vector nohost bind("string") +#pragma acc routine (function) vector nohost bind("string") // CHECK: #pragma acc routine(function) device_type(Something) seq #pragma acc routine(function) device_type(Something) seq @@ -82,20 +82,20 @@ struct DepS { #pragma acc routine (MemFunc) device_type(Lambda) vector }; -// CHECK: #pragma acc routine(DepS<int>::Lambda) gang -#pragma acc routine(DepS<int>::Lambda) gang +// CHECK: #pragma acc routine(DepS<int>::Lambda) gang bind("string") +#pragma acc routine(DepS<int>::Lambda) gang bind("string") // CHECK: #pragma acc routine(DepS<int>::MemFunc) gang(dim: 1) #pragma acc routine(DepS<int>::MemFunc) gang(dim:1) -// CHECK: #pragma acc routine(DepS<int>::StaticMemFunc) vector -#pragma acc routine(DepS<int>::StaticMemFunc) vector +// CHECK: #pragma acc routine(DepS<int>::StaticMemFunc) vector bind(identifier) +#pragma acc routine(DepS<int>::StaticMemFunc) vector bind(identifier) template<typename T> void TemplFunc() { // CHECK: #pragma acc routine(T::MemFunc) gang(dim: T::SomethingElse()) #pragma acc routine(T::MemFunc) gang(dim:T::SomethingElse()) -// CHECK: #pragma acc routine(T::StaticMemFunc) worker nohost -#pragma acc routine(T::StaticMemFunc) worker nohost -// CHECK: #pragma acc routine(T::Lambda) nohost seq -#pragma acc routine(T::Lambda) nohost seq +// CHECK: #pragma acc routine(T::StaticMemFunc) worker nohost bind(identifier) +#pragma acc routine(T::StaticMemFunc) worker nohost bind(identifier) +// CHECK: #pragma acc routine(T::Lambda) nohost seq bind("string") +#pragma acc routine(T::Lambda) nohost seq bind("string") } diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c index 0964ae78216dd..5178138216773 100644 --- a/clang/test/ParserOpenACC/parse-clauses.c +++ b/clang/test/ParserOpenACC/parse-clauses.c @@ -1286,13 +1286,7 @@ void BCP1(); // expected-error@+1{{expected identifier or string literal}} #pragma acc routine(BCP1) seq bind() - // expected-warning@+2{{OpenACC clause 'bind' not yet implemented, clause ignored}} // expected-warning@+1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}} #pragma acc routine seq bind("ReductionClauseParsing") -void BCP2(); - // expected-warning@+1{{OpenACC clause 'bind' not yet implemented, clause ignored}} -#pragma acc routine(BCP1) seq bind(BCP2) - - // expected-error@+1{{use of undeclared identifier 'unknown_thing'}} #pragma acc routine(BCP1) seq bind(unknown_thing) diff --git a/clang/test/ParserOpenACC/parse-clauses.cpp b/clang/test/ParserOpenACC/parse-clauses.cpp index 03a0fd945e0b9..6c84018060943 100644 --- a/clang/test/ParserOpenACC/parse-clauses.cpp +++ b/clang/test/ParserOpenACC/parse-clauses.cpp @@ -49,52 +49,15 @@ void use() { templ<7, S>(); } -namespace NS { -void NSFunc(); - -class RecordTy { // #RecTy - static constexpr bool Value = false; // #VAL - void priv_mem_function(); // #PrivMemFun - public: - static constexpr bool ValuePub = true; - void mem_function(); -}; -template<typename T> -class TemplTy{}; -void function(); -} - - - // expected-warning@+1{{OpenACC clause 'bind' not yet implemented, clause ignored}} +// expected-error@+2{{expected ')'}} +// expected-note@+1{{to match this '('}} #pragma acc routine(use) seq bind(NS::NSFunc) - // expected-error@+2{{'RecordTy' does not refer to a value}} - // expected-note@#RecTy{{declared here}} -#pragma acc routine(use) seq bind(NS::RecordTy) - // expected-error@+3{{'Value' is a private member of 'NS::RecordTy'}} - // expected-note@#VAL{{implicitly declared private here}} - // expected-warning@+1{{OpenACC clause 'bind' not yet implemented, clause ignored}} -#pragma acc routine(use) seq bind(NS::RecordTy::Value) - // expected-warning@+1{{OpenACC clause 'bind' not yet implemented, clause ignored}} -#pragma acc routine(use) seq bind(NS::RecordTy::ValuePub) - // expected-warning@+1{{OpenACC clause 'bind' not yet implemented, clause ignored}} -#pragma acc routine(use) seq bind(NS::TemplTy<int>) - // expected-error@+1{{no member named 'unknown' in namespace 'NS'}} -#pragma acc routine(use) seq bind(NS::unknown<int>) - // expected-warning@+1{{OpenACC clause 'bind' not yet implemented, clause ignored}} -#pragma acc routine(use) seq bind(NS::function) - // expected-error@+3{{'priv_mem_function' is a private member of 'NS::RecordTy'}} - // expected-note@#PrivMemFun{{implicitly declared private here}} - // expected-warning@+1{{OpenACC clause 'bind' not yet implemented, clause ignored}} -#pragma acc routine(use) seq bind(NS::RecordTy::priv_mem_function) - // expected-warning@+1{{OpenACC clause 'bind' not yet implemented, clause ignored}} -#pragma acc routine(use) seq bind(NS::RecordTy::mem_function) // expected-error@+1{{string literal with user-defined suffix cannot be used here}} #pragma acc routine(use) seq bind("unknown udl"_UDL) - // 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}} + // expected-warning@+1{{encoding prefix 'u' on an unevaluated string literal has no effect}} #pragma acc routine(use) seq bind(u"16 bits") - // 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) seq bind(U"32 bits") +void another_func(); + // expected-warning@+1{{encoding prefix 'U' on an unevaluated string literal has no effect}} +#pragma acc routine(another_func) seq bind(U"32 bits") diff --git a/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c b/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c index 1cdc657df0460..dbba86a1ffc61 100644 --- a/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c +++ b/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c @@ -125,7 +125,7 @@ void uses() { for(unsigned i = 0; i < 5; ++i); #pragma acc parallel loop auto collapse(1) for(unsigned i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}} + // expected-error@+1{{OpenACC 'bind' clause is not valid on 'parallel loop' directive}} #pragma acc parallel loop auto bind(Var) for(unsigned i = 0; i < 5; ++i); #pragma acc parallel loop auto vector_length(1) @@ -242,7 +242,7 @@ void uses() { for(unsigned i = 0; i < 5; ++i); #pragma acc parallel loop collapse(1) auto for(unsigned i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}} + // expected-error@+1{{OpenACC 'bind' clause is not valid on 'parallel loop' directive}} #pragma acc parallel loop bind(Var) auto for(unsigned i = 0; i < 5; ++i); #pragma acc parallel loop vector_length(1) auto @@ -360,7 +360,7 @@ void uses() { for(unsigned i = 0; i < 5; ++i); #pragma acc parallel loop independent collapse(1) for(unsigned i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}} + // expected-error@+1{{OpenACC 'bind' clause is not valid on 'parallel loop' directive}} #pragma acc parallel loop independent bind(Var) for(unsigned i = 0; i < 5; ++i); #pragma acc parallel loop independent vector_length(1) @@ -477,7 +477,7 @@ void uses() { for(unsigned i = 0; i < 5; ++i); #pragma acc parallel loop collapse(1) independent for(unsigned i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}} + // expected-error@+1{{OpenACC 'bind' clause is not valid on 'parallel loop' directive}} #pragma acc parallel loop bind(Var) independent for(unsigned i = 0; i < 5; ++i); #pragma acc parallel loop vector_length(1) independent @@ -603,7 +603,7 @@ void uses() { for(unsigned i = 0; i < 5; ++i); #pragma acc parallel loop seq collapse(1) for(unsigned i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}} + // expected-error@+1{{OpenACC 'bind' clause is not valid on 'parallel loop' directive}} #pragma acc parallel loop seq bind(Var) for(unsigned i = 0; i < 5; ++i); #pragma acc parallel loop seq vector_length(1) @@ -726,7 +726,7 @@ void uses() { for(unsigned i = 0; i < 5; ++i); #pragma acc parallel loop collapse(1) seq for(unsigned i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}} + // expected-error@+1{{OpenACC 'bind' clause is not valid on 'parallel loop' directive}} #pragma acc parallel loop bind(Var) seq for(unsigned i = 0; i < 5; ++i); #pragma acc parallel loop vector_length(1) seq diff --git a/clang/test/SemaOpenACC/combined-construct-device_type-clause.c b/clang/test/SemaOpenACC/combined-construct-device_type-clause.c index 49987663c79e1..38c72551997b8 100644 --- a/clang/test/SemaOpenACC/combined-construct-device_type-clause.c +++ b/clang/test/SemaOpenACC/combined-construct-device_type-clause.c @@ -173,8 +173,7 @@ void uses() { for(int i = 0; i < 5; ++i); #pragma acc serial loop device_type(*) collapse(1) for(int i = 0; i < 5; ++i); - // expected-error@+2{{OpenACC clause 'bind' may not follow a 'device_type' clause in a 'parallel loop' construct}} - // expected-note@+1{{previous clause is here}} + // expected-error@+1{{OpenACC 'bind' clause is not valid on 'parallel loop' directive}} #pragma acc parallel loop device_type(*) bind(Var) for(int i = 0; i < 5; ++i); // expected-error@+1{{OpenACC 'vector_length' clause is not valid on 'serial loop' directive}} diff --git a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c index a41936b57f09a..1e74e5ac23377 100644 --- a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c +++ b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c @@ -180,8 +180,7 @@ void uses() { // expected-error@+1{{OpenACC 'collapse' clause is not valid on 'kernels' directive}} #pragma acc kernels device_type(*) collapse(1) while(1); - // expected-error@+2{{OpenACC clause 'bind' may not follow a 'device_type' clause in a 'kernels' construct}} - // expected-note@+1{{previous clause is here}} + // expected-error@+1{{OpenACC 'bind' clause is not valid on 'kernels' directive}} #pragma acc kernels device_type(*) bind(Var) while(1); #pragma acc kernels device_type(*) vector_length(1) diff --git a/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c b/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c index 9dde1a96720d0..2f2bd15935856 100644 --- a/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c +++ b/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c @@ -150,7 +150,7 @@ void uses() { for(unsigned i = 0; i < 5; ++i); #pragma acc loop auto collapse(1) for(unsigned i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}} + // expected-error@+1{{OpenACC 'bind' clause is not valid on 'loop' directive}} #pragma acc loop auto bind(Var) for(unsigned i = 0; i < 5; ++i); // expected-error@+1{{OpenACC 'vector_length' clause is not valid on 'loop' directive}} @@ -284,7 +284,7 @@ void uses() { for(unsigned i = 0; i < 5; ++i); #pragma acc loop collapse(1) auto for(unsigned i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}} + // expected-error@+1{{OpenACC 'bind' clause is not valid on 'loop' directive}} #pragma acc loop bind(Var) auto for(unsigned i = 0; i < 5; ++i); // expected-error@+1{{OpenACC 'vector_length' clause is not valid on 'loop' directive}} @@ -419,7 +419,7 @@ void uses() { for(unsigned i = 0; i < 5; ++i); #pragma acc loop independent collapse(1) for(unsigned i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}} + // expected-error@+1{{OpenACC 'bind' clause is not valid on 'loop' directive}} #pragma acc loop independent bind(Var) for(unsigned i = 0; i < 5; ++i); // expected-error@+1{{OpenACC 'vector_length' clause is not valid on 'loop' directive}} @@ -553,7 +553,7 @@ void uses() { for(unsigned i = 0; i < 5; ++i); #pragma acc loop collapse(1) independent for(unsigned i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}} + // expected-error@+1{{OpenACC 'bind' clause is not valid on 'loop' directive}} #pragma acc loop bind(Var) independent for(unsigned i = 0; i < 5; ++i); // expected-error@+1{{OpenACC 'vector_length' clause is not valid on 'loop' directive}} @@ -696,7 +696,7 @@ void uses() { for(unsigned i = 0; i < 5; ++i); #pragma acc loop seq collapse(1) for(unsigned i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}} + // expected-error@+1{{OpenACC 'bind' clause is not valid on 'loop' directive}} #pragma acc loop seq bind(Var) for(unsigned i = 0; i < 5; ++i); // expected-error@+1{{OpenACC 'vector_length' clause is not valid on 'loop' directive}} @@ -836,7 +836,7 @@ void uses() { for(unsigned i = 0; i < 5; ++i); #pragma acc loop collapse(1) seq for(unsigned i = 0; i < 5; ++i); - // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}} + // expected-error@+1{{OpenACC 'bind' clause is not valid on 'loop' directive}} #pragma acc loop bind(Var) seq for(unsigned i = 0; i < 5; ++i); // expected-error@+1{{OpenACC 'vector_length' clause is not valid on 'loop' directive}} diff --git a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c index 83100b80bff91..5c5e76881a889 100644 --- a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c +++ b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c @@ -152,8 +152,7 @@ void uses() { for(int i = 0; i < 5; ++i); #pragma acc loop device_type(*) collapse(1) for(int i = 0; i < 5; ++i); - // expected-error@+2{{OpenACC clause 'bind' may not follow a 'device_type' clause in a 'loop' construct}} - // expected-note@+1{{previous clause is here}} + // expected-error@+1{{OpenACC 'bind' clause is not valid on 'loop' directive}} #pragma acc loop device_type(*) bind(Var) for(int i = 0; i < 5; ++i); // expected-error@+1{{OpenACC 'vector_length' clause is not valid on 'loop' directive}} diff --git a/clang/test/SemaOpenACC/routine-construct-ast.cpp b/clang/test/SemaOpenACC/routine-construct-ast.cpp index f9510eae31590..7715ec62e7d4e 100644 --- a/clang/test/SemaOpenACC/routine-construct-ast.cpp +++ b/clang/test/SemaOpenACC/routine-construct-ast.cpp @@ -7,17 +7,20 @@ #ifndef PCH_HELPER #define PCH_HELPER auto Lambda = [](){}; -#pragma acc routine(Lambda) worker nohost +#pragma acc routine(Lambda) worker nohost bind("string") // CHECK: OpenACCRoutineDecl{{.*}} routine name_specified // CHECK-NEXT: DeclRefExpr{{.*}} 'Lambda' '(lambda at // CHECK-NEXT: worker clause // CHECK-NEXT: nohost clause +// CHECK-NEXT: bind clause +// CHECK-NEXT: StringLiteral{{.*}} "string" int function(); -#pragma acc routine (function) nohost vector +#pragma acc routine (function) nohost vector bind(identifier) // CHECK: OpenACCRoutineDecl{{.*}} routine name_specified // CHECK-NEXT: DeclRefExpr{{.*}} 'function' 'int ()' // CHECK-NEXT: nohost clause // CHECK-NEXT: vector clause +// CHECK-NEXT: bind clause identifier 'identifier' #pragma acc routine(function) device_type(Something) seq // CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified @@ -263,20 +266,23 @@ void TemplFunc() { // CHECK-NEXT: NestedNameSpecifier{{.*}} 'T' // CHECK-NEXT: gang clause // CHECK-NEXT: CallExpr{{.*}}'<dependent type>' -#pragma acc routine(T::StaticMemFunc) nohost worker +#pragma acc routine(T::StaticMemFunc) nohost worker bind("string") // CHECK-NEXT: DeclStmt // CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}}'<dependent type>' // CHECK-NEXT: NestedNameSpecifier{{.*}} 'T' // CHECK-NEXT: nohost clause // CHECK-NEXT: worker clause -#pragma acc routine(T::Lambda) seq nohost +// CHECK-NEXT: bind clause +// CHECK-NEXT: StringLiteral{{.*}} "string" +#pragma acc routine(T::Lambda) seq nohost bind(identifier) // CHECK-NEXT: DeclStmt // CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}}'<dependent type>' // CHECK-NEXT: NestedNameSpecifier{{.*}} 'T' // CHECK-NEXT: seq clause // CHECK-NEXT: nohost clause +// CHECK-NEXT: bind clause identifier 'identifier' // Instantiation: // CHECK: FunctionDecl{{.*}} TemplFunc 'void ()' implicit_instantiation @@ -294,6 +300,8 @@ void TemplFunc() { // CHECK-NEXT: NestedNameSpecifier{{.*}} 'S' // CHECK-NEXT: nohost clause // CHECK-NEXT: worker clause +// CHECK-NEXT: bind clause +// CHECK-NEXT: StringLiteral{{.*}} "string" // CHECK-NEXT: DeclStmt // CHECK-NEXT: OpenACCRoutineDecl{{.*}} routine name_specified @@ -301,6 +309,7 @@ void TemplFunc() { // CHECK-NEXT: NestedNameSpecifier{{.*}} 'S' // CHECK-NEXT: seq clause // CHECK-NEXT: nohost clause +// CHECK-NEXT: bind clause identifier 'identifier' } void usage() { diff --git a/clang/test/SemaOpenACC/routine-construct-clauses.cpp b/clang/test/SemaOpenACC/routine-construct-clauses.cpp index 5fea3d734a410..da8ed382c4d63 100644 --- a/clang/test/SemaOpenACC/routine-construct-clauses.cpp +++ b/clang/test/SemaOpenACC/routine-construct-clauses.cpp @@ -1,6 +1,7 @@ // RUN: %clang_cc1 %s -fopenacc -verify void Func(); +void Func2(); #pragma acc routine(Func) worker #pragma acc routine(Func) vector nohost @@ -131,10 +132,8 @@ void Inst() { #pragma acc routine(Func) device_type(*) worker #pragma acc routine(Func) device_type(*) vector #pragma acc routine(Func) dtype(*) seq -// expected-warning@+1{{OpenACC clause 'bind' not yet implemented}} #pragma acc routine(Func) seq device_type(*) bind("asdf") -// expected-warning@+1{{OpenACC clause 'bind' not yet implemented}} -#pragma acc routine(Func) seq device_type(*) bind(getSomeInt) +#pragma acc routine(Func2) seq device_type(*) bind(WhateverElse) #pragma acc routine(Func) seq dtype(*) device_type(*) // expected-error@+2{{OpenACC clause 'nohost' may not follow a 'dtype' clause in a 'routine' construct}} // expected-note@+1{{previous clause is here}} @@ -142,3 +141,27 @@ void Inst() { // expected-error@+2{{OpenACC clause 'nohost' may not follow a 'device_type' clause in a 'routine' construct}} // expected-note@+1{{previous clause is here}} #pragma acc routine(Func) seq device_type(*) nohost + +// 2.15: a bind clause may not bind to a routine name that has a visible bind clause. +void Func3(); +#pragma acc routine(Func3) seq bind("asdf") +// OK: Doesn't have a bind +#pragma acc routine(Func3) seq +// expected-error@+2{{multiple 'routine' directives with 'bind' clauses are not permitted to refer to the same function}} +// expected-note@-4{{previous clause is here}} +#pragma acc routine(Func3) seq bind("asdf") + +void Func4(); +#pragma acc routine(Func4) seq +// OK: Doesn't have a bind +#pragma acc routine(Func4) seq bind("asdf") +// expected-error@+2{{multiple 'routine' directives with 'bind' clauses are not permitted to refer to the same function}} +// expected-note@-2{{previous clause is here}} +#pragma acc routine(Func4) seq bind("asdf") +void Func5(); +#pragma acc routine(Func5) seq bind("asdf") +// expected-error@+2{{multiple 'routine' directives with 'bind' clauses are not permitted to refer to the same function}} +// expected-note@-2{{previous clause is here}} +#pragma acc routine(Func5) seq bind("asdf") +// OK: Doesn't have a bind +#pragma acc routine(Func5) seq diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp index f412a38a92337..fa7aec175e294 100644 --- a/clang/tools/libclang/CIndex.cpp +++ b/clang/tools/libclang/CIndex.cpp @@ -2974,6 +2974,9 @@ void OpenACCClauseEnqueue::VisitIndependentClause( const OpenACCIndependentClause &C) {} void OpenACCClauseEnqueue::VisitSeqClause(const OpenACCSeqClause &C) {} void OpenACCClauseEnqueue::VisitNoHostClause(const OpenACCNoHostClause &C) {} +void OpenACCClauseEnqueue::VisitBindClause(const OpenACCBindClause &C) { + assert(false && "TODO ERICH"); +} void OpenACCClauseEnqueue::VisitFinalizeClause(const OpenACCFinalizeClause &C) { } void OpenACCClauseEnqueue::VisitIfPresentClause( _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
