[clang] 8a8f135 - [OpenACC] Implement 'bind' ast/sema for 'routine' directive
via cfe-commits
cfe-commits at lists.llvm.org
Mon Mar 10 07:49:21 PDT 2025
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 at +1{{expected identifier or string literal}}
#pragma acc routine(BCP1) seq bind()
- // expected-warning at +2{{OpenACC clause 'bind' not yet implemented, clause ignored}}
// expected-warning at +1{{OpenACC construct 'routine' with implicit function not yet implemented, pragma ignored}}
#pragma acc routine seq bind("ReductionClauseParsing")
-void BCP2();
- // expected-warning at +1{{OpenACC clause 'bind' not yet implemented, clause ignored}}
-#pragma acc routine(BCP1) seq bind(BCP2)
-
- // expected-error at +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 at +1{{OpenACC clause 'bind' not yet implemented, clause ignored}}
+// expected-error at +2{{expected ')'}}
+// expected-note at +1{{to match this '('}}
#pragma acc routine(use) seq bind(NS::NSFunc)
- // expected-error at +2{{'RecordTy' does not refer to a value}}
- // expected-note@#RecTy{{declared here}}
-#pragma acc routine(use) seq bind(NS::RecordTy)
- // expected-error at +3{{'Value' is a private member of 'NS::RecordTy'}}
- // expected-note@#VAL{{implicitly declared private here}}
- // expected-warning at +1{{OpenACC clause 'bind' not yet implemented, clause ignored}}
-#pragma acc routine(use) seq bind(NS::RecordTy::Value)
- // expected-warning at +1{{OpenACC clause 'bind' not yet implemented, clause ignored}}
-#pragma acc routine(use) seq bind(NS::RecordTy::ValuePub)
- // expected-warning at +1{{OpenACC clause 'bind' not yet implemented, clause ignored}}
-#pragma acc routine(use) seq bind(NS::TemplTy<int>)
- // expected-error at +1{{no member named 'unknown' in namespace 'NS'}}
-#pragma acc routine(use) seq bind(NS::unknown<int>)
- // expected-warning at +1{{OpenACC clause 'bind' not yet implemented, clause ignored}}
-#pragma acc routine(use) seq bind(NS::function)
- // expected-error at +3{{'priv_mem_function' is a private member of 'NS::RecordTy'}}
- // expected-note@#PrivMemFun{{implicitly declared private here}}
- // expected-warning at +1{{OpenACC clause 'bind' not yet implemented, clause ignored}}
-#pragma acc routine(use) seq bind(NS::RecordTy::priv_mem_function)
- // expected-warning at +1{{OpenACC clause 'bind' not yet implemented, clause ignored}}
-#pragma acc routine(use) seq bind(NS::RecordTy::mem_function)
// expected-error at +1{{string literal with user-defined suffix cannot be used here}}
#pragma acc routine(use) seq bind("unknown udl"_UDL)
- // expected-warning at +2{{encoding prefix 'u' on an unevaluated string literal has no effect}}
- // expected-warning at +1{{OpenACC clause 'bind' not yet implemented, clause ignored}}
+ // expected-warning at +1{{encoding prefix 'u' on an unevaluated string literal has no effect}}
#pragma acc routine(use) seq bind(u"16 bits")
- // expected-warning at +2{{encoding prefix 'U' on an unevaluated string literal has no effect}}
- // expected-warning at +1{{OpenACC clause 'bind' not yet implemented, clause ignored}}
-#pragma acc routine(use) seq bind(U"32 bits")
+void another_func();
+ // expected-warning at +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 at +1{{OpenACC clause 'bind' not yet implemented}}
+ // expected-error at +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 at +1{{OpenACC clause 'bind' not yet implemented}}
+ // expected-error at +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 at +1{{OpenACC clause 'bind' not yet implemented}}
+ // expected-error at +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 at +1{{OpenACC clause 'bind' not yet implemented}}
+ // expected-error at +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 at +1{{OpenACC clause 'bind' not yet implemented}}
+ // expected-error at +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 at +1{{OpenACC clause 'bind' not yet implemented}}
+ // expected-error at +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 at +2{{OpenACC clause 'bind' may not follow a 'device_type' clause in a 'parallel loop' construct}}
- // expected-note at +1{{previous clause is here}}
+ // expected-error at +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 at +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 at +1{{OpenACC 'collapse' clause is not valid on 'kernels' directive}}
#pragma acc kernels device_type(*) collapse(1)
while(1);
- // expected-error at +2{{OpenACC clause 'bind' may not follow a 'device_type' clause in a 'kernels' construct}}
- // expected-note at +1{{previous clause is here}}
+ // expected-error at +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 at +1{{OpenACC clause 'bind' not yet implemented}}
+ // expected-error at +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 at +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 at +1{{OpenACC clause 'bind' not yet implemented}}
+ // expected-error at +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 at +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 at +1{{OpenACC clause 'bind' not yet implemented}}
+ // expected-error at +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 at +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 at +1{{OpenACC clause 'bind' not yet implemented}}
+ // expected-error at +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 at +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 at +1{{OpenACC clause 'bind' not yet implemented}}
+ // expected-error at +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 at +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 at +1{{OpenACC clause 'bind' not yet implemented}}
+ // expected-error at +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 at +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 at +2{{OpenACC clause 'bind' may not follow a 'device_type' clause in a 'loop' construct}}
- // expected-note at +1{{previous clause is here}}
+ // expected-error at +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 at +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 at +1{{OpenACC clause 'bind' not yet implemented}}
#pragma acc routine(Func) seq device_type(*) bind("asdf")
-// expected-warning at +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 at +2{{OpenACC clause 'nohost' may not follow a 'dtype' clause in a 'routine' construct}}
// expected-note at +1{{previous clause is here}}
@@ -142,3 +141,27 @@ void Inst() {
// expected-error at +2{{OpenACC clause 'nohost' may not follow a 'device_type' clause in a 'routine' construct}}
// expected-note at +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 at +2{{multiple 'routine' directives with 'bind' clauses are not permitted to refer to the same function}}
+// expected-note at -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 at +2{{multiple 'routine' directives with 'bind' clauses are not permitted to refer to the same function}}
+// expected-note at -2{{previous clause is here}}
+#pragma acc routine(Func4) seq bind("asdf")
+void Func5();
+#pragma acc routine(Func5) seq bind("asdf")
+// expected-error at +2{{multiple 'routine' directives with 'bind' clauses are not permitted to refer to the same function}}
+// expected-note at -2{{previous clause is here}}
+#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(
More information about the cfe-commits
mailing list