[clang] 48c8a57 - [OpenACC] Implement 'deviceptr' and 'attach' sema for compute constructs
via cfe-commits
cfe-commits at lists.llvm.org
Mon May 6 09:29:09 PDT 2024
Author: erichkeane
Date: 2024-05-06T09:29:04-07:00
New Revision: 48c8a5791ae71c96661479f684459b7b9427a22d
URL: https://github.com/llvm/llvm-project/commit/48c8a5791ae71c96661479f684459b7b9427a22d
DIFF: https://github.com/llvm/llvm-project/commit/48c8a5791ae71c96661479f684459b7b9427a22d.diff
LOG: [OpenACC] Implement 'deviceptr' and 'attach' sema for compute constructs
These two are very similar to the other 'var-list' variants, except they
require that the type of the variable be a pointer. This patch
implements that restriction.
Added:
clang/test/SemaOpenACC/compute-construct-attach-clause.c
clang/test/SemaOpenACC/compute-construct-attach-clause.cpp
clang/test/SemaOpenACC/compute-construct-deviceptr-clause.c
clang/test/SemaOpenACC/compute-construct-deviceptr-clause.cpp
Modified:
clang/include/clang/AST/OpenACCClause.h
clang/include/clang/Basic/DiagnosticSemaKinds.td
clang/include/clang/Basic/OpenACCClauses.def
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/TreeTransform.h
clang/lib/Serialization/ASTReader.cpp
clang/lib/Serialization/ASTWriter.cpp
clang/test/AST/ast-print-openacc-compute-construct.cpp
clang/test/ParserOpenACC/parse-clauses.c
clang/test/SemaOpenACC/compute-construct-varlist-ast.cpp
clang/tools/libclang/CIndex.cpp
Removed:
################################################################################
diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h
index abbd7c3989bcca..ec6b4aebcb9f4c 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -313,6 +313,44 @@ class OpenACCFirstPrivateClause final
ArrayRef<Expr *> VarList, SourceLocation EndLoc);
};
+class OpenACCDevicePtrClause final
+ : public OpenACCClauseWithVarList,
+ public llvm::TrailingObjects<OpenACCDevicePtrClause, Expr *> {
+
+ OpenACCDevicePtrClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
+ ArrayRef<Expr *> VarList, SourceLocation EndLoc)
+ : OpenACCClauseWithVarList(OpenACCClauseKind::DevicePtr, BeginLoc,
+ LParenLoc, EndLoc) {
+ std::uninitialized_copy(VarList.begin(), VarList.end(),
+ getTrailingObjects<Expr *>());
+ setExprs(MutableArrayRef(getTrailingObjects<Expr *>(), VarList.size()));
+ }
+
+public:
+ static OpenACCDevicePtrClause *
+ Create(const ASTContext &C, SourceLocation BeginLoc, SourceLocation LParenLoc,
+ ArrayRef<Expr *> VarList, SourceLocation EndLoc);
+};
+
+class OpenACCAttachClause final
+ : public OpenACCClauseWithVarList,
+ public llvm::TrailingObjects<OpenACCAttachClause, Expr *> {
+
+ OpenACCAttachClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
+ ArrayRef<Expr *> VarList, SourceLocation EndLoc)
+ : OpenACCClauseWithVarList(OpenACCClauseKind::Attach, BeginLoc, LParenLoc,
+ EndLoc) {
+ std::uninitialized_copy(VarList.begin(), VarList.end(),
+ getTrailingObjects<Expr *>());
+ setExprs(MutableArrayRef(getTrailingObjects<Expr *>(), VarList.size()));
+ }
+
+public:
+ static OpenACCAttachClause *
+ Create(const ASTContext &C, SourceLocation BeginLoc, SourceLocation LParenLoc,
+ ArrayRef<Expr *> VarList, SourceLocation EndLoc);
+};
+
class OpenACCNoCreateClause final
: public OpenACCClauseWithVarList,
public llvm::TrailingObjects<OpenACCNoCreateClause, Expr *> {
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 6c4d92790afc51..9a0bae9c216de9 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12337,4 +12337,7 @@ def warn_acc_deprecated_alias_name
: Warning<"OpenACC clause name '%0' is a deprecated clause name and is "
"now an alias for '%1'">,
InGroup<DiagGroup<"openacc-deprecated-clause-alias">>;
+def err_acc_var_not_pointer_type
+ : Error<"expected pointer in '%0' clause, type is %1">;
+def note_acc_expected_pointer_var : Note<"expected variable of pointer type">;
} // end of sema component.
diff --git a/clang/include/clang/Basic/OpenACCClauses.def b/clang/include/clang/Basic/OpenACCClauses.def
index 06c7a379b22261..c92e5eb1e1b634 100644
--- a/clang/include/clang/Basic/OpenACCClauses.def
+++ b/clang/include/clang/Basic/OpenACCClauses.def
@@ -21,6 +21,7 @@
#define CLAUSE_ALIAS(ALIAS_NAME, CLAUSE_NAME)
#endif
+VISIT_CLAUSE(Attach)
VISIT_CLAUSE(Copy)
CLAUSE_ALIAS(PCopy, Copy)
CLAUSE_ALIAS(PresentOrCopy, Copy)
@@ -34,6 +35,7 @@ VISIT_CLAUSE(Create)
CLAUSE_ALIAS(PCreate, Create)
CLAUSE_ALIAS(PresentOrCreate, Create)
VISIT_CLAUSE(Default)
+VISIT_CLAUSE(DevicePtr)
VISIT_CLAUSE(FirstPrivate)
VISIT_CLAUSE(If)
VISIT_CLAUSE(NoCreate)
diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h
index 5f77ec90d0b650..32d94ee8f33fed 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -134,6 +134,8 @@ class SemaOpenACC : public SemaBase {
ClauseKind == OpenACCClauseKind::Create ||
ClauseKind == OpenACCClauseKind::PCreate ||
ClauseKind == OpenACCClauseKind::PresentOrCreate ||
+ ClauseKind == OpenACCClauseKind::Attach ||
+ ClauseKind == OpenACCClauseKind::DevicePtr ||
ClauseKind == OpenACCClauseKind::FirstPrivate) &&
"Parsed clause kind does not have a var-list");
return std::get<VarListDetails>(Details).VarList;
@@ -217,6 +219,8 @@ class SemaOpenACC : public SemaBase {
ClauseKind == OpenACCClauseKind::Create ||
ClauseKind == OpenACCClauseKind::PCreate ||
ClauseKind == OpenACCClauseKind::PresentOrCreate ||
+ ClauseKind == OpenACCClauseKind::Attach ||
+ ClauseKind == OpenACCClauseKind::DevicePtr ||
ClauseKind == OpenACCClauseKind::FirstPrivate) &&
"Parsed clause kind does not have a var-list");
assert((!IsReadOnly || ClauseKind == OpenACCClauseKind::CopyIn ||
@@ -251,6 +255,8 @@ class SemaOpenACC : public SemaBase {
ClauseKind == OpenACCClauseKind::Create ||
ClauseKind == OpenACCClauseKind::PCreate ||
ClauseKind == OpenACCClauseKind::PresentOrCreate ||
+ ClauseKind == OpenACCClauseKind::Attach ||
+ ClauseKind == OpenACCClauseKind::DevicePtr ||
ClauseKind == OpenACCClauseKind::FirstPrivate) &&
"Parsed clause kind does not have a var-list");
assert((!IsReadOnly || ClauseKind == OpenACCClauseKind::CopyIn ||
@@ -315,6 +321,10 @@ class SemaOpenACC : public SemaBase {
/// declaration reference to a variable of the correct type.
ExprResult ActOnVar(Expr *VarExpr);
+ /// Called to check the 'var' type is a variable of pointer type, necessary
+ /// for 'deviceptr' and 'attach' clauses. Returns true on success.
+ bool CheckVarIsPointerType(OpenACCClauseKind ClauseKind, Expr *VarExpr);
+
/// Checks and creates an Array Section used in an OpenACC construct/clause.
ExprResult ActOnArraySectionExpr(Expr *Base, SourceLocation LBLoc,
Expr *LowerBound,
diff --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp
index f682100e35d37b..c1affa97b781ca 100644
--- a/clang/lib/AST/OpenACCClause.cpp
+++ b/clang/lib/AST/OpenACCClause.cpp
@@ -156,6 +156,26 @@ OpenACCFirstPrivateClause *OpenACCFirstPrivateClause::Create(
OpenACCFirstPrivateClause(BeginLoc, LParenLoc, VarList, EndLoc);
}
+OpenACCAttachClause *OpenACCAttachClause::Create(const ASTContext &C,
+ SourceLocation BeginLoc,
+ SourceLocation LParenLoc,
+ ArrayRef<Expr *> VarList,
+ SourceLocation EndLoc) {
+ void *Mem =
+ C.Allocate(OpenACCAttachClause::totalSizeToAlloc<Expr *>(VarList.size()));
+ return new (Mem) OpenACCAttachClause(BeginLoc, LParenLoc, VarList, EndLoc);
+}
+
+OpenACCDevicePtrClause *OpenACCDevicePtrClause::Create(const ASTContext &C,
+ SourceLocation BeginLoc,
+ SourceLocation LParenLoc,
+ ArrayRef<Expr *> VarList,
+ SourceLocation EndLoc) {
+ void *Mem = C.Allocate(
+ OpenACCDevicePtrClause::totalSizeToAlloc<Expr *>(VarList.size()));
+ return new (Mem) OpenACCDevicePtrClause(BeginLoc, LParenLoc, VarList, EndLoc);
+}
+
OpenACCNoCreateClause *OpenACCNoCreateClause::Create(const ASTContext &C,
SourceLocation BeginLoc,
SourceLocation LParenLoc,
@@ -282,6 +302,21 @@ void OpenACCClausePrinter::VisitFirstPrivateClause(
OS << ")";
}
+void OpenACCClausePrinter::VisitAttachClause(const OpenACCAttachClause &C) {
+ OS << "attach(";
+ llvm::interleaveComma(C.getVarList(), OS,
+ [&](const Expr *E) { printExpr(E); });
+ OS << ")";
+}
+
+void OpenACCClausePrinter::VisitDevicePtrClause(
+ const OpenACCDevicePtrClause &C) {
+ OS << "deviceptr(";
+ llvm::interleaveComma(C.getVarList(), OS,
+ [&](const Expr *E) { printExpr(E); });
+ OS << ")";
+}
+
void OpenACCClausePrinter::VisitNoCreateClause(const OpenACCNoCreateClause &C) {
OS << "no_create(";
llvm::interleaveComma(C.getVarList(), OS,
diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index b97c351f83dbf4..11d3f3d4cec444 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -2543,6 +2543,18 @@ void OpenACCClauseProfiler::VisitFirstPrivateClause(
Profiler.VisitStmt(E);
}
+void OpenACCClauseProfiler::VisitAttachClause(
+ const OpenACCAttachClause &Clause) {
+ for (auto *E : Clause.getVarList())
+ Profiler.VisitStmt(E);
+}
+
+void OpenACCClauseProfiler::VisitDevicePtrClause(
+ const OpenACCDevicePtrClause &Clause) {
+ for (auto *E : Clause.getVarList())
+ Profiler.VisitStmt(E);
+}
+
void OpenACCClauseProfiler::VisitNoCreateClause(
const OpenACCNoCreateClause &Clause) {
for (auto *E : Clause.getVarList())
diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp
index 5aeeb80e8fe710..21167ca56e5947 100644
--- a/clang/lib/AST/TextNodeDumper.cpp
+++ b/clang/lib/AST/TextNodeDumper.cpp
@@ -397,10 +397,12 @@ void TextNodeDumper::Visit(const OpenACCClause *C) {
case OpenACCClauseKind::Default:
OS << '(' << cast<OpenACCDefaultClause>(C)->getDefaultClauseKind() << ')';
break;
+ case OpenACCClauseKind::Attach:
case OpenACCClauseKind::Copy:
case OpenACCClauseKind::PCopy:
case OpenACCClauseKind::PresentOrCopy:
case OpenACCClauseKind::If:
+ case OpenACCClauseKind::DevicePtr:
case OpenACCClauseKind::FirstPrivate:
case OpenACCClauseKind::NoCreate:
case OpenACCClauseKind::NumGangs:
diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp
index c90ae8806f0868..b4b81e2ba13ea6 100644
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@@ -945,17 +945,21 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
// make sure we get the right
diff erentiator.
assert(DirKind == OpenACCDirectiveKind::Update);
[[fallthrough]];
- case OpenACCClauseKind::Attach:
case OpenACCClauseKind::Delete:
case OpenACCClauseKind::Detach:
case OpenACCClauseKind::Device:
case OpenACCClauseKind::DeviceResident:
- case OpenACCClauseKind::DevicePtr:
case OpenACCClauseKind::Host:
case OpenACCClauseKind::Link:
case OpenACCClauseKind::UseDevice:
ParseOpenACCVarList();
break;
+ case OpenACCClauseKind::Attach:
+ case OpenACCClauseKind::DevicePtr:
+ // TODO: ERICH: Figure out how to limit to just ptrs?
+ ParsedClause.setVarListDetails(ParseOpenACCVarList(),
+ /*IsReadOnly=*/false, /*IsZero=*/false);
+ break;
case OpenACCClauseKind::Copy:
case OpenACCClauseKind::PCopy:
case OpenACCClauseKind::PresentOrCopy:
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 4e1a8c1277c5f0..8cf829cf215b54 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -170,6 +170,34 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
default:
return false;
}
+ case OpenACCClauseKind::Attach:
+ switch (DirectiveKind) {
+ case OpenACCDirectiveKind::Parallel:
+ case OpenACCDirectiveKind::Serial:
+ case OpenACCDirectiveKind::Kernels:
+ case OpenACCDirectiveKind::Data:
+ case OpenACCDirectiveKind::EnterData:
+ case OpenACCDirectiveKind::ParallelLoop:
+ case OpenACCDirectiveKind::SerialLoop:
+ case OpenACCDirectiveKind::KernelsLoop:
+ return true;
+ default:
+ return false;
+ }
+ case OpenACCClauseKind::DevicePtr:
+ switch (DirectiveKind) {
+ case OpenACCDirectiveKind::Parallel:
+ case OpenACCDirectiveKind::Serial:
+ case OpenACCDirectiveKind::Kernels:
+ case OpenACCDirectiveKind::Data:
+ case OpenACCDirectiveKind::Declare:
+ case OpenACCDirectiveKind::ParallelLoop:
+ case OpenACCDirectiveKind::SerialLoop:
+ case OpenACCDirectiveKind::KernelsLoop:
+ return true;
+ default:
+ return false;
+ }
default:
// Do nothing so we can go to the 'unimplemented' diagnostic instead.
return true;
@@ -513,6 +541,48 @@ SemaOpenACC::ActOnClause(ArrayRef<const OpenACCClause *> ExistingClauses,
Clause.getLParenLoc(), Clause.isZero(),
Clause.getVarList(), Clause.getEndLoc());
}
+ case OpenACCClauseKind::Attach: {
+ // Restrictions only properly implemented on 'compute' constructs, and
+ // 'compute' constructs are the only construct that can do anything with
+ // this yet, so skip/treat as unimplemented in this case.
+ if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()))
+ break;
+
+ // ActOnVar ensured that everything is a valid variable reference, but we
+ // still have to make sure it is a pointer type.
+ llvm::SmallVector<Expr *> VarList{Clause.getVarList().begin(),
+ Clause.getVarList().end()};
+ VarList.erase(std::remove_if(VarList.begin(), VarList.end(), [&](Expr *E) {
+ return CheckVarIsPointerType(OpenACCClauseKind::Attach, E);
+ }), VarList.end());
+ Clause.setVarListDetails(VarList,
+ /*IsReadOnly=*/false, /*IsZero=*/false);
+
+ return OpenACCAttachClause::Create(getASTContext(), Clause.getBeginLoc(),
+ Clause.getLParenLoc(),
+ Clause.getVarList(), Clause.getEndLoc());
+ }
+ case OpenACCClauseKind::DevicePtr: {
+ // Restrictions only properly implemented on 'compute' constructs, and
+ // 'compute' constructs are the only construct that can do anything with
+ // this yet, so skip/treat as unimplemented in this case.
+ if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()))
+ break;
+
+ // ActOnVar ensured that everything is a valid variable reference, but we
+ // still have to make sure it is a pointer type.
+ llvm::SmallVector<Expr *> VarList{Clause.getVarList().begin(),
+ Clause.getVarList().end()};
+ VarList.erase(std::remove_if(VarList.begin(), VarList.end(), [&](Expr *E) {
+ return CheckVarIsPointerType(OpenACCClauseKind::DevicePtr, E);
+ }), VarList.end());
+ Clause.setVarListDetails(VarList,
+ /*IsReadOnly=*/false, /*IsZero=*/false);
+
+ return OpenACCDevicePtrClause::Create(
+ getASTContext(), Clause.getBeginLoc(), Clause.getLParenLoc(),
+ Clause.getVarList(), Clause.getEndLoc());
+ }
default:
break;
}
@@ -641,6 +711,36 @@ ExprResult SemaOpenACC::ActOnIntExpr(OpenACCDirectiveKind DK,
return IntExpr;
}
+bool SemaOpenACC::CheckVarIsPointerType(OpenACCClauseKind ClauseKind,
+ Expr *VarExpr) {
+ // We already know that VarExpr is a proper reference to a variable, so we
+ // should be able to just take the type of the expression to get the type of
+ // the referenced variable.
+
+ // We've already seen an error, don't diagnose anything else.
+ if (!VarExpr || VarExpr->containsErrors())
+ return false;
+
+ if (isa<ArraySectionExpr>(VarExpr->IgnoreParenImpCasts()) ||
+ VarExpr->hasPlaceholderType(BuiltinType::ArraySection)) {
+ Diag(VarExpr->getExprLoc(), diag::err_array_section_use) << /*OpenACC=*/0;
+ Diag(VarExpr->getExprLoc(), diag::note_acc_expected_pointer_var);
+ return true;
+ }
+
+ QualType Ty = VarExpr->getType();
+ Ty = Ty.getNonReferenceType().getUnqualifiedType();
+
+ // Nothing we can do if this is a dependent type.
+ if (Ty->isDependentType())
+ return false;
+
+ if (!Ty->isPointerType())
+ return Diag(VarExpr->getExprLoc(), diag::err_acc_var_not_pointer_type)
+ << ClauseKind << Ty;
+ return false;
+}
+
ExprResult SemaOpenACC::ActOnVar(Expr *VarExpr) {
// We still need to retain the array subscript/subarray exprs, so work on a
// copy.
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index edf55703874688..a4ca8b5771a9f2 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -11318,6 +11318,43 @@ void OpenACCClauseTransform<Derived>::VisitCreateClause(
ParsedClause.isZero(), ParsedClause.getVarList(),
ParsedClause.getEndLoc());
}
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitAttachClause(
+ const OpenACCAttachClause &C) {
+ llvm::SmallVector<Expr *> VarList = VisitVarList(C.getVarList());
+
+ // Ensure each var is a pointer type.
+ VarList.erase(std::remove_if(VarList.begin(), VarList.end(), [&](Expr *E) {
+ return Self.getSema().OpenACC().CheckVarIsPointerType(
+ OpenACCClauseKind::Attach, E);
+ }), VarList.end());
+
+ ParsedClause.setVarListDetails(VarList,
+ /*IsReadOnly=*/false, /*IsZero=*/false);
+ NewClause = OpenACCAttachClause::Create(
+ Self.getSema().getASTContext(), ParsedClause.getBeginLoc(),
+ ParsedClause.getLParenLoc(), ParsedClause.getVarList(),
+ ParsedClause.getEndLoc());
+}
+
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitDevicePtrClause(
+ const OpenACCDevicePtrClause &C) {
+ llvm::SmallVector<Expr *> VarList = VisitVarList(C.getVarList());
+
+ // Ensure each var is a pointer type.
+ VarList.erase(std::remove_if(VarList.begin(), VarList.end(), [&](Expr *E) {
+ return Self.getSema().OpenACC().CheckVarIsPointerType(
+ OpenACCClauseKind::DevicePtr, E);
+ }), VarList.end());
+
+ ParsedClause.setVarListDetails(VarList,
+ /*IsReadOnly=*/false, /*IsZero=*/false);
+ NewClause = OpenACCDevicePtrClause::Create(
+ Self.getSema().getASTContext(), ParsedClause.getBeginLoc(),
+ ParsedClause.getLParenLoc(), ParsedClause.getVarList(),
+ ParsedClause.getEndLoc());
+}
template <typename Derived>
void OpenACCClauseTransform<Derived>::VisitNumWorkersClause(
diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
index 4050d66d45170c..81b78edd9c6cfe 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -11822,6 +11822,18 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
return OpenACCFirstPrivateClause::Create(getContext(), BeginLoc, LParenLoc,
VarList, EndLoc);
}
+ case OpenACCClauseKind::Attach: {
+ SourceLocation LParenLoc = readSourceLocation();
+ llvm::SmallVector<Expr *> VarList = readOpenACCVarList();
+ return OpenACCAttachClause::Create(getContext(), BeginLoc, LParenLoc,
+ VarList, EndLoc);
+ }
+ case OpenACCClauseKind::DevicePtr: {
+ SourceLocation LParenLoc = readSourceLocation();
+ llvm::SmallVector<Expr *> VarList = readOpenACCVarList();
+ return OpenACCDevicePtrClause::Create(getContext(), BeginLoc, LParenLoc,
+ VarList, EndLoc);
+ }
case OpenACCClauseKind::NoCreate: {
SourceLocation LParenLoc = readSourceLocation();
llvm::SmallVector<Expr *> VarList = readOpenACCVarList();
@@ -11879,11 +11891,9 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
case OpenACCClauseKind::Vector:
case OpenACCClauseKind::NoHost:
case OpenACCClauseKind::UseDevice:
- case OpenACCClauseKind::Attach:
case OpenACCClauseKind::Delete:
case OpenACCClauseKind::Detach:
case OpenACCClauseKind::Device:
- case OpenACCClauseKind::DevicePtr:
case OpenACCClauseKind::DeviceResident:
case OpenACCClauseKind::Host:
case OpenACCClauseKind::Link:
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index cf77b4c0df2bbb..8a0116fa893247 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -7820,6 +7820,18 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
writeOpenACCVarList(FPC);
return;
}
+ case OpenACCClauseKind::Attach: {
+ const auto *AC = cast<OpenACCAttachClause>(C);
+ writeSourceLocation(AC->getLParenLoc());
+ writeOpenACCVarList(AC);
+ return;
+ }
+ case OpenACCClauseKind::DevicePtr: {
+ const auto *DPC = cast<OpenACCDevicePtrClause>(C);
+ writeSourceLocation(DPC->getLParenLoc());
+ writeOpenACCVarList(DPC);
+ return;
+ }
case OpenACCClauseKind::NoCreate: {
const auto *NCC = cast<OpenACCNoCreateClause>(C);
writeSourceLocation(NCC->getLParenLoc());
@@ -7877,11 +7889,9 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
case OpenACCClauseKind::Vector:
case OpenACCClauseKind::NoHost:
case OpenACCClauseKind::UseDevice:
- case OpenACCClauseKind::Attach:
case OpenACCClauseKind::Delete:
case OpenACCClauseKind::Detach:
case OpenACCClauseKind::Device:
- case OpenACCClauseKind::DevicePtr:
case OpenACCClauseKind::DeviceResident:
case OpenACCClauseKind::Host:
case OpenACCClauseKind::Link:
diff --git a/clang/test/AST/ast-print-openacc-compute-construct.cpp b/clang/test/AST/ast-print-openacc-compute-construct.cpp
index 112f328f5cb9ce..1ee1e15bdfc3ac 100644
--- a/clang/test/AST/ast-print-openacc-compute-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-compute-construct.cpp
@@ -2,7 +2,9 @@
void foo() {
int i;
+ int *iPtr;
float array[5];
+ float *arrayPtr[5];
// CHECK: #pragma acc parallel default(none)
// CHECK-NEXT: while (true)
#pragma acc parallel default(none)
@@ -65,5 +67,13 @@ void foo() {
// CHECK: #pragma acc parallel create(i, array[1], array, array[1:2]) pcreate(zero: i, array[1], array, array[1:2]) present_or_create(i, array[1], array, array[1:2])
#pragma acc parallel create(i, array[1], array, array[1:2]) pcreate(zero: i, array[1], array, array[1:2]) present_or_create(i, array[1], array, array[1:2])
while(true);
+
+ // CHECK: #pragma acc serial attach(iPtr, arrayPtr[0])
+#pragma acc serial attach(iPtr, arrayPtr[0])
+ while(true);
+
+ // CHECK: #pragma acc kernels deviceptr(iPtr, arrayPtr[0])
+#pragma acc kernels deviceptr(iPtr, arrayPtr[0])
+ while(true);
}
diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c
index 65247e4db63efa..035b7ab4c1f40f 100644
--- a/clang/test/ParserOpenACC/parse-clauses.c
+++ b/clang/test/ParserOpenACC/parse-clauses.c
@@ -548,26 +548,30 @@ void VarListClauses() {
#pragma acc serial present(s.array[s.value : 5], s.value), seq
for(;;){}
- // expected-error at +3{{expected ','}}
- // expected-warning at +2{{OpenACC clause 'deviceptr' not yet implemented, clause ignored}}
+
+ void *IsPointer;
+ // expected-error at +5{{expected ','}}
+ // expected-error at +4{{expected pointer in 'deviceptr' clause, type is 'char'}}
+ // expected-error at +3{{OpenACC sub-array is not allowed here}}
+ // expected-note at +2{{expected variable of pointer type}}
// expected-warning at +1{{OpenACC clause 'seq' not yet implemented, clause ignored}}
#pragma acc serial deviceptr(s.array[s.value] s.array[s.value :5] ), seq
for(;;){}
- // expected-warning at +2{{OpenACC clause 'deviceptr' not yet implemented, clause ignored}}
// expected-warning at +1{{OpenACC clause 'seq' not yet implemented, clause ignored}}
-#pragma acc serial deviceptr(s.array[s.value : 5], s.value), seq
+#pragma acc serial deviceptr(IsPointer), seq
for(;;){}
- // expected-error at +3{{expected ','}}
- // expected-warning at +2{{OpenACC clause 'attach' not yet implemented, clause ignored}}
+ // expected-error at +5{{expected ','}}
+ // expected-error at +4{{expected pointer in 'attach' clause, type is 'char'}}
+ // expected-error at +3{{OpenACC sub-array is not allowed here}}
+ // expected-note at +2{{expected variable of pointer type}}
// expected-warning at +1{{OpenACC clause 'seq' not yet implemented, clause ignored}}
#pragma acc serial attach(s.array[s.value] s.array[s.value :5] ), seq
for(;;){}
- // expected-warning at +2{{OpenACC clause 'attach' not yet implemented, clause ignored}}
// expected-warning at +1{{OpenACC clause 'seq' not yet implemented, clause ignored}}
-#pragma acc serial attach(s.array[s.value : 5], s.value), seq
+#pragma acc serial attach(IsPointer), seq
for(;;){}
// expected-error at +3{{expected ','}}
diff --git a/clang/test/SemaOpenACC/compute-construct-attach-clause.c b/clang/test/SemaOpenACC/compute-construct-attach-clause.c
new file mode 100644
index 00000000000000..de735308528adb
--- /dev/null
+++ b/clang/test/SemaOpenACC/compute-construct-attach-clause.c
@@ -0,0 +1,61 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+struct S {
+ int IntMem;
+ int *PtrMem;
+};
+
+void uses() {
+ int LocalInt;
+ int *LocalPtr;
+ int Array[5];
+ int *PtrArray[5];
+ struct S s;
+
+ // expected-error at +1{{expected pointer in 'attach' clause, type is 'int'}}
+#pragma acc parallel attach(LocalInt)
+ while (1);
+
+ // expected-error at +1{{OpenACC variable is not a valid variable name, sub-array, array element, or composite variable member}}
+#pragma acc parallel attach(&LocalInt)
+ while (1);
+
+#pragma acc serial attach(LocalPtr)
+ while (1);
+
+ // expected-error at +1{{expected pointer in 'attach' clause, type is 'int[5]'}}
+#pragma acc kernels attach(Array)
+ while (1);
+
+ // expected-error at +1{{expected pointer in 'attach' clause, type is 'int'}}
+#pragma acc parallel attach(Array[0])
+ while (1);
+
+ // expected-error at +2{{OpenACC sub-array is not allowed here}}
+ // expected-note at +1{{expected variable of pointer type}}
+#pragma acc parallel attach(Array[0:1])
+ while (1);
+
+ // expected-error at +1{{expected pointer in 'attach' clause, type is 'int *[5]'}}
+#pragma acc parallel attach(PtrArray)
+ while (1);
+
+#pragma acc parallel attach(PtrArray[0])
+ while (1);
+
+ // expected-error at +2{{OpenACC sub-array is not allowed here}}
+ // expected-note at +1{{expected variable of pointer type}}
+#pragma acc parallel attach(PtrArray[0:1])
+ while (1);
+
+ // expected-error at +1{{expected pointer in 'attach' clause, type is 'struct S'}}
+#pragma acc parallel attach(s)
+ while (1);
+
+ // expected-error at +1{{expected pointer in 'attach' clause, type is 'int'}}
+#pragma acc parallel attach(s.IntMem)
+ while (1);
+
+#pragma acc parallel attach(s.PtrMem)
+ while (1);
+}
diff --git a/clang/test/SemaOpenACC/compute-construct-attach-clause.cpp b/clang/test/SemaOpenACC/compute-construct-attach-clause.cpp
new file mode 100644
index 00000000000000..a89d346c2645ab
--- /dev/null
+++ b/clang/test/SemaOpenACC/compute-construct-attach-clause.cpp
@@ -0,0 +1,120 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+struct S {
+ int IntMem;
+ int *PtrMem;
+ operator int*();
+};
+
+void uses() {
+ int LocalInt;
+ int *LocalPtr;
+ int Array[5];
+ int *PtrArray[5];
+ struct S s;
+
+ // expected-error at +1{{expected pointer in 'attach' clause, type is 'int'}}
+#pragma acc parallel attach(LocalInt)
+ while (true);
+
+#pragma acc parallel attach(LocalPtr)
+ while (true);
+
+ // expected-error at +1{{expected pointer in 'attach' clause, type is 'int[5]'}}
+#pragma acc parallel attach(Array)
+ while (true);
+
+ // expected-error at +1{{expected pointer in 'attach' clause, type is 'int'}}
+#pragma acc parallel attach(Array[0])
+ while (true);
+
+ // expected-error at +2{{OpenACC sub-array is not allowed here}}
+ // expected-note at +1{{expected variable of pointer type}}
+#pragma acc parallel attach(Array[0:1])
+ while (true);
+
+ // expected-error at +1{{expected pointer in 'attach' clause, type is 'int *[5]'}}
+#pragma acc parallel attach(PtrArray)
+ while (true);
+
+#pragma acc parallel attach(PtrArray[0])
+ while (true);
+
+ // expected-error at +2{{OpenACC sub-array is not allowed here}}
+ // expected-note at +1{{expected variable of pointer type}}
+#pragma acc parallel attach(PtrArray[0:1])
+ while (true);
+
+ // expected-error at +1{{expected pointer in 'attach' clause, type is 'struct S'}}
+#pragma acc parallel attach(s)
+ while (true);
+
+ // expected-error at +1{{expected pointer in 'attach' clause, type is 'int'}}
+#pragma acc parallel attach(s.IntMem)
+ while (true);
+
+#pragma acc parallel attach(s.PtrMem)
+ while (true);
+}
+
+template<typename T, typename TPtr, typename TStruct, auto &R1>
+void Templ() {
+ T SomeInt;
+ TPtr SomePtr;
+ T SomeIntArray[5];
+ TPtr SomeIntPtrArray[5];
+ TStruct SomeStruct;
+
+ // expected-error at +2{{expected pointer in 'attach' clause, type is 'int'}}
+ // expected-note@#INST{{in instantiation of function template specialization}}
+#pragma acc parallel attach(SomeInt)
+ while (true);
+
+#pragma acc parallel attach(SomePtr)
+ while (true);
+
+ // expected-error at +1{{expected pointer in 'attach' clause, type is 'int[5]'}}
+#pragma acc parallel attach(SomeIntArray)
+ while (true);
+
+ // expected-error at +1{{expected pointer in 'attach' clause, type is 'int'}}
+#pragma acc parallel attach(SomeIntArray[0])
+ while (true);
+
+ // expected-error at +2{{OpenACC sub-array is not allowed here}}
+ // expected-note at +1{{expected variable of pointer type}}
+#pragma acc parallel attach(SomeIntArray[0:1])
+ while (true);
+
+ // expected-error at +1{{expected pointer in 'attach' clause, type is 'int *[5]'}}
+#pragma acc parallel attach(SomeIntPtrArray)
+ while (true);
+
+#pragma acc parallel attach(SomeIntPtrArray[0])
+ while (true);
+
+ // expected-error at +2{{OpenACC sub-array is not allowed here}}
+ // expected-note at +1{{expected variable of pointer type}}
+#pragma acc parallel attach(SomeIntPtrArray[0:1])
+ while (true);
+
+ // expected-error at +1{{expected pointer in 'attach' clause, type is 'S'}}
+#pragma acc parallel attach(SomeStruct)
+ while (true);
+
+ // expected-error at +1{{expected pointer in 'attach' clause, type is 'int'}}
+#pragma acc parallel attach(SomeStruct.IntMem)
+ while (true);
+
+#pragma acc parallel attach(SomeStruct.PtrMem)
+ while (true);
+
+ // expected-error at +1{{expected pointer in 'attach' clause, type is 'int'}}
+#pragma acc parallel attach(R1)
+ while (true);
+}
+
+void inst() {
+ static constexpr int CEVar = 1;
+ Templ<int, int*, S, CEVar>(); // #INST
+}
diff --git a/clang/test/SemaOpenACC/compute-construct-deviceptr-clause.c b/clang/test/SemaOpenACC/compute-construct-deviceptr-clause.c
new file mode 100644
index 00000000000000..e5d328eb0b28bc
--- /dev/null
+++ b/clang/test/SemaOpenACC/compute-construct-deviceptr-clause.c
@@ -0,0 +1,61 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+struct S {
+ int IntMem;
+ int *PtrMem;
+};
+
+void uses() {
+ int LocalInt;
+ int *LocalPtr;
+ int Array[5];
+ int *PtrArray[5];
+ struct S s;
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'int'}}
+#pragma acc parallel deviceptr(LocalInt)
+ while (1);
+
+ // expected-error at +1{{OpenACC variable is not a valid variable name, sub-array, array element, or composite variable member}}
+#pragma acc parallel deviceptr(&LocalInt)
+ while (1);
+
+#pragma acc serial deviceptr(LocalPtr)
+ while (1);
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'int[5]'}}
+#pragma acc kernels deviceptr(Array)
+ while (1);
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'int'}}
+#pragma acc parallel deviceptr(Array[0])
+ while (1);
+
+ // expected-error at +2{{OpenACC sub-array is not allowed here}}
+ // expected-note at +1{{expected variable of pointer type}}
+#pragma acc parallel deviceptr(Array[0:1])
+ while (1);
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'int *[5]'}}
+#pragma acc parallel deviceptr(PtrArray)
+ while (1);
+
+#pragma acc parallel deviceptr(PtrArray[0])
+ while (1);
+
+ // expected-error at +2{{OpenACC sub-array is not allowed here}}
+ // expected-note at +1{{expected variable of pointer type}}
+#pragma acc parallel deviceptr(PtrArray[0:1])
+ while (1);
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'struct S'}}
+#pragma acc parallel deviceptr(s)
+ while (1);
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'int'}}
+#pragma acc parallel deviceptr(s.IntMem)
+ while (1);
+
+#pragma acc parallel deviceptr(s.PtrMem)
+ while (1);
+}
diff --git a/clang/test/SemaOpenACC/compute-construct-deviceptr-clause.cpp b/clang/test/SemaOpenACC/compute-construct-deviceptr-clause.cpp
new file mode 100644
index 00000000000000..83409c91d4818f
--- /dev/null
+++ b/clang/test/SemaOpenACC/compute-construct-deviceptr-clause.cpp
@@ -0,0 +1,120 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+struct S {
+ int IntMem;
+ int *PtrMem;
+ operator int*();
+};
+
+void uses() {
+ int LocalInt;
+ int *LocalPtr;
+ int Array[5];
+ int *PtrArray[5];
+ struct S s;
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'int'}}
+#pragma acc parallel deviceptr(LocalInt)
+ while (true);
+
+#pragma acc parallel deviceptr(LocalPtr)
+ while (true);
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'int[5]'}}
+#pragma acc parallel deviceptr(Array)
+ while (true);
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'int'}}
+#pragma acc parallel deviceptr(Array[0])
+ while (true);
+
+ // expected-error at +2{{OpenACC sub-array is not allowed here}}
+ // expected-note at +1{{expected variable of pointer type}}
+#pragma acc parallel deviceptr(Array[0:1])
+ while (true);
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'int *[5]'}}
+#pragma acc parallel deviceptr(PtrArray)
+ while (true);
+
+#pragma acc parallel deviceptr(PtrArray[0])
+ while (true);
+
+ // expected-error at +2{{OpenACC sub-array is not allowed here}}
+ // expected-note at +1{{expected variable of pointer type}}
+#pragma acc parallel deviceptr(PtrArray[0:1])
+ while (true);
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'struct S'}}
+#pragma acc parallel deviceptr(s)
+ while (true);
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'int'}}
+#pragma acc parallel deviceptr(s.IntMem)
+ while (true);
+
+#pragma acc parallel deviceptr(s.PtrMem)
+ while (true);
+}
+
+template<typename T, typename TPtr, typename TStruct, auto &R1>
+void Templ() {
+ T SomeInt;
+ TPtr SomePtr;
+ T SomeIntArray[5];
+ TPtr SomeIntPtrArray[5];
+ TStruct SomeStruct;
+
+ // expected-error at +2{{expected pointer in 'deviceptr' clause, type is 'int'}}
+ // expected-note@#INST{{in instantiation of function template specialization}}
+#pragma acc parallel deviceptr(SomeInt)
+ while (true);
+
+#pragma acc parallel deviceptr(SomePtr)
+ while (true);
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'int[5]'}}
+#pragma acc parallel deviceptr(SomeIntArray)
+ while (true);
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'int'}}
+#pragma acc parallel deviceptr(SomeIntArray[0])
+ while (true);
+
+ // expected-error at +2{{OpenACC sub-array is not allowed here}}
+ // expected-note at +1{{expected variable of pointer type}}
+#pragma acc parallel deviceptr(SomeIntArray[0:1])
+ while (true);
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'int *[5]'}}
+#pragma acc parallel deviceptr(SomeIntPtrArray)
+ while (true);
+
+#pragma acc parallel deviceptr(SomeIntPtrArray[0])
+ while (true);
+
+ // expected-error at +2{{OpenACC sub-array is not allowed here}}
+ // expected-note at +1{{expected variable of pointer type}}
+#pragma acc parallel deviceptr(SomeIntPtrArray[0:1])
+ while (true);
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'S'}}
+#pragma acc parallel deviceptr(SomeStruct)
+ while (true);
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'int'}}
+#pragma acc parallel deviceptr(SomeStruct.IntMem)
+ while (true);
+
+#pragma acc parallel deviceptr(SomeStruct.PtrMem)
+ while (true);
+
+ // expected-error at +1{{expected pointer in 'deviceptr' clause, type is 'int'}}
+#pragma acc parallel deviceptr(R1)
+ while (true);
+}
+
+void inst() {
+ static constexpr int CEVar = 1;
+ Templ<int, int*, S, CEVar>(); // #INST
+}
diff --git a/clang/test/SemaOpenACC/compute-construct-varlist-ast.cpp b/clang/test/SemaOpenACC/compute-construct-varlist-ast.cpp
index d35f62adfe0794..e057678d924957 100644
--- a/clang/test/SemaOpenACC/compute-construct-varlist-ast.cpp
+++ b/clang/test/SemaOpenACC/compute-construct-varlist-ast.cpp
@@ -191,6 +191,17 @@ void NormalUses(float *PointerParam) {
// CHECK-NEXT: WhileStmt
// CHECK-NEXT: CXXBoolLiteralExpr
// CHECK-NEXT: NullStmt
+
+#pragma acc parallel attach(PointerParam) deviceptr(PointerParam)
+ while (true);
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+ // CHECK-NEXT: attach clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'float *' lvalue ParmVar{{.*}} 'PointerParam' 'float *'
+ // CHECK-NEXT: deviceptr clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'float *' lvalue ParmVar{{.*}} 'PointerParam' 'float *'
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: NullStmt
}
// This example is an error typically, but we want to make sure we're properly
@@ -402,6 +413,17 @@ void TemplUses(T t, U u, T*PointerParam) {
// CHECK-NEXT: CXXBoolLiteralExpr
// CHECK-NEXT: NullStmt
+#pragma acc parallel attach(PointerParam) deviceptr(PointerParam)
+ while (true);
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+ // CHECK-NEXT: attach clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'T *' lvalue ParmVar{{.*}} 'PointerParam' 'T *'
+ // CHECK-NEXT: deviceptr clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'T *' lvalue ParmVar{{.*}} 'PointerParam' 'T *'
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: NullStmt
+
// CHECK-NEXT: DeclStmt
// CHECK-NEXT: VarDecl{{.*}}EndMarker
int EndMarker;
@@ -604,6 +626,16 @@ void TemplUses(T t, U u, T*PointerParam) {
// CHECK-NEXT: CXXBoolLiteralExpr
// CHECK-NEXT: NullStmt
+//#pragma acc parallel attach(PointerParam) deviceptr(PointerParam)
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+ // CHECK-NEXT: attach clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'int *' lvalue ParmVar{{.*}} 'PointerParam' 'int *'
+ // CHECK-NEXT: deviceptr clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'int *' lvalue ParmVar{{.*}} 'PointerParam' 'int *'
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: NullStmt
+
// CHECK-NEXT: DeclStmt
// CHECK-NEXT: VarDecl{{.*}}EndMarker
}
@@ -613,6 +645,8 @@ struct S {
// CHECK: CXXRecordDecl{{.*}} implicit struct S
int ThisMember;
// CHECK-NEXT: FieldDecl{{.*}} ThisMember 'int'
+ int *ThisMemberPtr;
+ // CHECK-NEXT: FieldDecl{{.*}} ThisMemberPtr 'int *'
int ThisMemberArray[5];
// CHECK-NEXT: FieldDecl{{.*}} ThisMemberArray 'int[5]'
@@ -620,10 +654,11 @@ struct S {
// CHECK-NEXT: CXXMethodDecl{{.*}} foo 'void ()'
template<typename T>
- void bar() {
+ void bar(T *PointerParam) {
// CHECK-NEXT: FunctionTemplateDecl{{.*}}bar
// CHECK-NEXT: TemplateTypeParmDecl{{.*}}typename depth 0 index 0 T
- // CHECK-NEXT: CXXMethodDecl{{.*}} bar 'void ()' implicit-inline
+ // CHECK-NEXT: CXXMethodDecl{{.*}} bar 'void (T *)' implicit-inline
+ // CHECK-NEXT: ParmVarDecl{{.*}} PointerParam 'T *'
// CHECK-NEXT: CompoundStmt
#pragma acc parallel private(ThisMember, this->ThisMemberArray[1])
@@ -664,10 +699,28 @@ struct S {
// CHECK-NEXT: CXXBoolLiteralExpr
// CHECK-NEXT: NullStmt
+#pragma acc parallel attach(PointerParam, this, this->ThisMemberPtr) deviceptr(PointerParam, this, ThisMemberPtr)
+ while (true);
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+ // CHECK-NEXT: attach clause
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'T *' lvalue ParmVar{{.*}} 'PointerParam' 'T *'
+ // CHECK-NEXT: CXXThisExpr{{.*}} 'S *' this
+ // CHECK-NEXT: MemberExpr{{.*}} 'int *' lvalue ->ThisMemberPtr
+ // CHECK-NEXT: CXXThisExpr{{.*}} 'S *' this
+ // CHECK-NEXT: deviceptr clause
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'T *' lvalue ParmVar{{.*}} 'PointerParam' 'T *'
+ // CHECK-NEXT: CXXThisExpr{{.*}} 'S *' this
+ // CHECK-NEXT: MemberExpr{{.*}} 'int *' lvalue ->ThisMemberPtr
+ // CHECK-NEXT: CXXThisExpr{{.*}} 'S *' implicit this
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: NullStmt
+
// Check Instantiations:
- // CHECK-NEXT: CXXMethodDecl{{.*}} used bar 'void ()' implicit_instantiation implicit-inline
+ // CHECK-NEXT: CXXMethodDecl{{.*}} used bar 'void (int *)' implicit_instantiation implicit-inline
// CHECK-NEXT: TemplateArgument type 'int'
// CHECK-NEXT: BuiltinType{{.*}} 'int'
+ // CHECK-NEXT: ParmVarDecl{{.*}} PointerParam 'int *'
// CHECK-NEXT: CompoundStmt
// #pragma acc parallel private(ThisMember, this->ThisMemberArray[1])
@@ -704,6 +757,22 @@ struct S {
// CHECK-NEXT: WhileStmt
// CHECK-NEXT: CXXBoolLiteralExpr
// CHECK-NEXT: NullStmt
+
+//#pragma acc parallel attach(PointerParam, this, this->ThisMemberPtr) deviceptr(PointerParam, this, ThisMemberPtr)
+ // CHECK-NEXT: OpenACCComputeConstruct{{.*}} parallel
+ // CHECK-NEXT: attach clause
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'int *' lvalue ParmVar{{.*}} 'PointerParam' 'int *'
+ // CHECK-NEXT: CXXThisExpr{{.*}} 'S *' this
+ // CHECK-NEXT: MemberExpr{{.*}} 'int *' lvalue ->ThisMemberPtr
+ // CHECK-NEXT: CXXThisExpr{{.*}} 'S *' this
+ // CHECK-NEXT: deviceptr clause
+ // CHECK-NEXT: DeclRefExpr{{.*}} 'int *' lvalue ParmVar{{.*}} 'PointerParam' 'int *'
+ // CHECK-NEXT: CXXThisExpr{{.*}} 'S *' this
+ // CHECK-NEXT: MemberExpr{{.*}} 'int *' lvalue ->ThisMemberPtr
+ // CHECK-NEXT: CXXThisExpr{{.*}} 'S *' implicit this
+ // CHECK-NEXT: WhileStmt
+ // CHECK-NEXT: CXXBoolLiteralExpr
+ // CHECK-NEXT: NullStmt
}
};
@@ -906,7 +975,7 @@ void Inst() {
TemplUses<CEVar, int, int[1]>({}, {}, &i);
S s;
- s.bar<int>();
+ s.bar<int>(&i);
STempl<int> stempl;
stempl.bar<int>();
}
diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
index 487f5785957685..6c07c4d2e30738 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -2840,6 +2840,13 @@ void OpenACCClauseEnqueue::VisitCopyOutClause(const OpenACCCopyOutClause &C) {
void OpenACCClauseEnqueue::VisitCreateClause(const OpenACCCreateClause &C) {
VisitVarList(C);
}
+void OpenACCClauseEnqueue::VisitAttachClause(const OpenACCAttachClause &C) {
+ VisitVarList(C);
+}
+void OpenACCClauseEnqueue::VisitDevicePtrClause(
+ const OpenACCDevicePtrClause &C) {
+ VisitVarList(C);
+}
} // namespace
void EnqueueVisitor::EnqueueChildren(const OpenACCClause *C) {
More information about the cfe-commits
mailing list