[clang] 3351b3b - [OpenACC] implement 'detach' clause sema
via cfe-commits
cfe-commits at lists.llvm.org
Fri Dec 13 13:51:47 PST 2024
Author: erichkeane
Date: 2024-12-13T13:51:41-08:00
New Revision: 3351b3bf8dcb9aebfa6f491fcbe5a00acbcc3291
URL: https://github.com/llvm/llvm-project/commit/3351b3bf8dcb9aebfa6f491fcbe5a00acbcc3291
DIFF: https://github.com/llvm/llvm-project/commit/3351b3bf8dcb9aebfa6f491fcbe5a00acbcc3291.diff
LOG: [OpenACC] implement 'detach' clause sema
This is another new clause specific to 'exit data' that takes a pointer
argument. This patch implements this the same way we do a few other
clauses (like attach) that have the same restrictions.
Added:
clang/test/SemaOpenACC/data-construct-detach-ast.cpp
clang/test/SemaOpenACC/data-construct-detach-clause.c
Modified:
clang/include/clang/AST/OpenACCClause.h
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-data-construct.cpp
clang/test/ParserOpenACC/parse-clauses.c
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/data-construct.cpp
clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
clang/test/SemaOpenACC/loop-construct-device_type-clause.c
clang/tools/libclang/CIndex.cpp
Removed:
################################################################################
diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h
index 57cd9a1e8564d2..7778f8199b3af6 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -744,6 +744,28 @@ class OpenACCAttachClause final
ArrayRef<Expr *> VarList, SourceLocation EndLoc);
};
+class OpenACCDetachClause final
+ : public OpenACCClauseWithVarList,
+ public llvm::TrailingObjects<OpenACCDetachClause, Expr *> {
+
+ OpenACCDetachClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
+ ArrayRef<Expr *> VarList, SourceLocation EndLoc)
+ : OpenACCClauseWithVarList(OpenACCClauseKind::Detach, BeginLoc, LParenLoc,
+ EndLoc) {
+ std::uninitialized_copy(VarList.begin(), VarList.end(),
+ getTrailingObjects<Expr *>());
+ setExprs(MutableArrayRef(getTrailingObjects<Expr *>(), VarList.size()));
+ }
+
+public:
+ static bool classof(const OpenACCClause *C) {
+ return C->getClauseKind() == OpenACCClauseKind::Detach;
+ }
+ static OpenACCDetachClause *
+ 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/OpenACCClauses.def b/clang/include/clang/Basic/OpenACCClauses.def
index cc788e4de45279..87983c3849480f 100644
--- a/clang/include/clang/Basic/OpenACCClauses.def
+++ b/clang/include/clang/Basic/OpenACCClauses.def
@@ -38,6 +38,7 @@ VISIT_CLAUSE(Create)
CLAUSE_ALIAS(PCreate, Create, true)
CLAUSE_ALIAS(PresentOrCreate, Create, true)
VISIT_CLAUSE(Default)
+VISIT_CLAUSE(Detach)
VISIT_CLAUSE(DevicePtr)
VISIT_CLAUSE(DeviceType)
CLAUSE_ALIAS(DType, DeviceType, false)
diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h
index 170a6f4885c965..ea3f34e3f4a959 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -399,6 +399,7 @@ class SemaOpenACC : public SemaBase {
ClauseKind == OpenACCClauseKind::PCreate ||
ClauseKind == OpenACCClauseKind::PresentOrCreate ||
ClauseKind == OpenACCClauseKind::Attach ||
+ ClauseKind == OpenACCClauseKind::Detach ||
ClauseKind == OpenACCClauseKind::DevicePtr ||
ClauseKind == OpenACCClauseKind::Reduction ||
ClauseKind == OpenACCClauseKind::FirstPrivate) &&
@@ -535,6 +536,7 @@ class SemaOpenACC : public SemaBase {
ClauseKind == OpenACCClauseKind::PCreate ||
ClauseKind == OpenACCClauseKind::PresentOrCreate ||
ClauseKind == OpenACCClauseKind::Attach ||
+ ClauseKind == OpenACCClauseKind::Detach ||
ClauseKind == OpenACCClauseKind::DevicePtr ||
ClauseKind == OpenACCClauseKind::FirstPrivate) &&
"Parsed clause kind does not have a var-list");
@@ -571,6 +573,7 @@ class SemaOpenACC : public SemaBase {
ClauseKind == OpenACCClauseKind::PCreate ||
ClauseKind == OpenACCClauseKind::PresentOrCreate ||
ClauseKind == OpenACCClauseKind::Attach ||
+ ClauseKind == OpenACCClauseKind::Detach ||
ClauseKind == OpenACCClauseKind::DevicePtr ||
ClauseKind == OpenACCClauseKind::FirstPrivate) &&
"Parsed clause kind does not have a var-list");
diff --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp
index 443cfa84474bed..d2d8f34e9014de 100644
--- a/clang/lib/AST/OpenACCClause.cpp
+++ b/clang/lib/AST/OpenACCClause.cpp
@@ -33,7 +33,8 @@ bool OpenACCClauseWithVarList::classof(const OpenACCClause *C) {
OpenACCFirstPrivateClause::classof(C) ||
OpenACCDevicePtrClause::classof(C) ||
OpenACCDevicePtrClause::classof(C) ||
- OpenACCAttachClause::classof(C) || OpenACCNoCreateClause::classof(C) ||
+ OpenACCDetachClause::classof(C) || OpenACCAttachClause::classof(C) ||
+ OpenACCNoCreateClause::classof(C) ||
OpenACCPresentClause::classof(C) || OpenACCCopyClause::classof(C) ||
OpenACCCopyInClause::classof(C) || OpenACCCopyOutClause::classof(C) ||
OpenACCReductionClause::classof(C) || OpenACCCreateClause::classof(C);
@@ -277,6 +278,16 @@ OpenACCAttachClause *OpenACCAttachClause::Create(const ASTContext &C,
return new (Mem) OpenACCAttachClause(BeginLoc, LParenLoc, VarList, EndLoc);
}
+OpenACCDetachClause *OpenACCDetachClause::Create(const ASTContext &C,
+ SourceLocation BeginLoc,
+ SourceLocation LParenLoc,
+ ArrayRef<Expr *> VarList,
+ SourceLocation EndLoc) {
+ void *Mem =
+ C.Allocate(OpenACCDetachClause::totalSizeToAlloc<Expr *>(VarList.size()));
+ return new (Mem) OpenACCDetachClause(BeginLoc, LParenLoc, VarList, EndLoc);
+}
+
OpenACCDevicePtrClause *OpenACCDevicePtrClause::Create(const ASTContext &C,
SourceLocation BeginLoc,
SourceLocation LParenLoc,
@@ -546,6 +557,13 @@ void OpenACCClausePrinter::VisitAttachClause(const OpenACCAttachClause &C) {
OS << ")";
}
+void OpenACCClausePrinter::VisitDetachClause(const OpenACCDetachClause &C) {
+ OS << "detach(";
+ llvm::interleaveComma(C.getVarList(), OS,
+ [&](const Expr *E) { printExpr(E); });
+ OS << ")";
+}
+
void OpenACCClausePrinter::VisitDevicePtrClause(
const OpenACCDevicePtrClause &C) {
OS << "deviceptr(";
diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index bd4956c15eea1c..4e9865f722f78e 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -2605,6 +2605,12 @@ void OpenACCClauseProfiler::VisitAttachClause(
Profiler.VisitStmt(E);
}
+void OpenACCClauseProfiler::VisitDetachClause(
+ const OpenACCDetachClause &Clause) {
+ for (auto *E : Clause.getVarList())
+ Profiler.VisitStmt(E);
+}
+
void OpenACCClauseProfiler::VisitDevicePtrClause(
const OpenACCDevicePtrClause &Clause) {
for (auto *E : Clause.getVarList())
diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp
index 0a5c0d561203af..b02b682fb0c58f 100644
--- a/clang/lib/AST/TextNodeDumper.cpp
+++ b/clang/lib/AST/TextNodeDumper.cpp
@@ -411,6 +411,7 @@ void TextNodeDumper::Visit(const OpenACCClause *C) {
case OpenACCClauseKind::If:
case OpenACCClauseKind::IfPresent:
case OpenACCClauseKind::Independent:
+ case OpenACCClauseKind::Detach:
case OpenACCClauseKind::DevicePtr:
case OpenACCClauseKind::Finalize:
case OpenACCClauseKind::FirstPrivate:
diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp
index 8c81936b35296c..5da7069edaa740 100644
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@@ -999,7 +999,6 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
assert(DirKind == OpenACCDirectiveKind::Update);
[[fallthrough]];
case OpenACCClauseKind::Delete:
- case OpenACCClauseKind::Detach:
case OpenACCClauseKind::Device:
case OpenACCClauseKind::DeviceResident:
case OpenACCClauseKind::Host:
@@ -1008,6 +1007,7 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
ParseOpenACCVarList(ClauseKind);
break;
case OpenACCClauseKind::Attach:
+ case OpenACCClauseKind::Detach:
case OpenACCClauseKind::DevicePtr:
ParsedClause.setVarListDetails(ParseOpenACCVarList(ClauseKind),
/*IsReadOnly=*/false, /*IsZero=*/false);
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 3917c9bf60b8c4..7156e37991284f 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -425,6 +425,14 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
return false;
}
}
+ case OpenACCClauseKind::Detach: {
+ switch (DirectiveKind) {
+ case OpenACCDirectiveKind::ExitData:
+ return true;
+ default:
+ return false;
+ }
+ }
}
default:
@@ -1043,6 +1051,21 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitAttachClause(
Clause.getEndLoc());
}
+OpenACCClause *SemaOpenACCClauseVisitor::VisitDetachClause(
+ SemaOpenACC::OpenACCParsedClause &Clause) {
+ // 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()};
+ llvm::erase_if(VarList, [&](Expr *E) {
+ return SemaRef.CheckVarIsPointerType(OpenACCClauseKind::Detach, E);
+ });
+ Clause.setVarListDetails(VarList,
+ /*IsReadOnly=*/false, /*IsZero=*/false);
+ return OpenACCDetachClause::Create(Ctx, Clause.getBeginLoc(),
+ Clause.getLParenLoc(), Clause.getVarList(),
+ Clause.getEndLoc());
+}
+
OpenACCClause *SemaOpenACCClauseVisitor::VisitDevicePtrClause(
SemaOpenACC::OpenACCParsedClause &Clause) {
// Restrictions only properly implemented on 'compute'/'combined'/'data'
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index e5c3584b4d8820..0f92186d08933c 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -11755,6 +11755,28 @@ void OpenACCClauseTransform<Derived>::VisitAttachClause(
ParsedClause.getEndLoc());
}
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitDetachClause(
+ const OpenACCDetachClause &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::Detach, E);
+ }),
+ VarList.end());
+
+ ParsedClause.setVarListDetails(VarList,
+ /*IsReadOnly=*/false, /*IsZero=*/false);
+ NewClause = OpenACCDetachClause::Create(
+ Self.getSema().getASTContext(), ParsedClause.getBeginLoc(),
+ ParsedClause.getLParenLoc(), ParsedClause.getVarList(),
+ ParsedClause.getEndLoc());
+}
+
template <typename Derived>
void OpenACCClauseTransform<Derived>::VisitDevicePtrClause(
const OpenACCDevicePtrClause &C) {
diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
index 4f7dc597d7bd5a..bed674233deecf 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -12442,6 +12442,12 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
return OpenACCAttachClause::Create(getContext(), BeginLoc, LParenLoc,
VarList, EndLoc);
}
+ case OpenACCClauseKind::Detach: {
+ SourceLocation LParenLoc = readSourceLocation();
+ llvm::SmallVector<Expr *> VarList = readOpenACCVarList();
+ return OpenACCDetachClause::Create(getContext(), BeginLoc, LParenLoc,
+ VarList, EndLoc);
+ }
case OpenACCClauseKind::DevicePtr: {
SourceLocation LParenLoc = readSourceLocation();
llvm::SmallVector<Expr *> VarList = readOpenACCVarList();
@@ -12586,7 +12592,6 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
case OpenACCClauseKind::NoHost:
case OpenACCClauseKind::UseDevice:
case OpenACCClauseKind::Delete:
- case OpenACCClauseKind::Detach:
case OpenACCClauseKind::Device:
case OpenACCClauseKind::DeviceResident:
case OpenACCClauseKind::Host:
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index f263082cc13e9a..0b2a62a219358c 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -8356,6 +8356,12 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
writeOpenACCVarList(AC);
return;
}
+ case OpenACCClauseKind::Detach: {
+ const auto *DC = cast<OpenACCDetachClause>(C);
+ writeSourceLocation(DC->getLParenLoc());
+ writeOpenACCVarList(DC);
+ return;
+ }
case OpenACCClauseKind::DevicePtr: {
const auto *DPC = cast<OpenACCDevicePtrClause>(C);
writeSourceLocation(DPC->getLParenLoc());
@@ -8501,7 +8507,6 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
case OpenACCClauseKind::NoHost:
case OpenACCClauseKind::UseDevice:
case OpenACCClauseKind::Delete:
- case OpenACCClauseKind::Detach:
case OpenACCClauseKind::Device:
case OpenACCClauseKind::DeviceResident:
case OpenACCClauseKind::Host:
diff --git a/clang/test/AST/ast-print-openacc-data-construct.cpp b/clang/test/AST/ast-print-openacc-data-construct.cpp
index bcc68a18232598..003dc03f342c15 100644
--- a/clang/test/AST/ast-print-openacc-data-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-data-construct.cpp
@@ -117,4 +117,7 @@ void foo() {
// CHECK: #pragma acc host_data if_present
#pragma acc host_data use_device(i) if_present
;
+// CHECK: #pragma acc exit data copyout(i) detach(iPtr, arrayPtr[0])
+#pragma acc exit data copyout(i) detach(iPtr, arrayPtr[0])
+
}
diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c
index 8003dad03c1c0d..0498fe16e512cd 100644
--- a/clang/test/ParserOpenACC/parse-clauses.c
+++ b/clang/test/ParserOpenACC/parse-clauses.c
@@ -488,14 +488,13 @@ void VarListClauses() {
#pragma acc serial attach(IsPointer), self
for(int i = 0; i < 5;++i) {}
- // expected-error at +2{{expected ','}}
- // expected-warning at +1{{OpenACC clause 'detach' not yet implemented, clause ignored}}
-#pragma acc serial detach(s.array[s.value] s.array[s.value :5] ), self
- for(int i = 0; i < 5;++i) {}
+ // expected-error at +4{{expected ','}}
+ // expected-error at +3{{expected pointer in 'detach' clause, type is 'char'}}
+ // expected-error at +2{{OpenACC sub-array is not allowed here}}
+ // expected-note at +1{{expected variable of pointer type}}
+#pragma acc exit data copyout(s) detach(s.array[s.value] s.array[s.value :5])
- // expected-warning at +1{{OpenACC clause 'detach' not yet implemented, clause ignored}}
-#pragma acc serial detach(s.array[s.value : 5], s.value), self
- for(int i = 0; i < 5;++i) {}
+#pragma acc exit data copyout(s) detach(IsPointer)
// expected-error at +1{{expected ','}}
#pragma acc serial private(s.array[s.value] s.array[s.value :5] ), self
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 516658f8e01ff4..b1d682c0ea5b39 100644
--- a/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
+++ b/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
@@ -72,7 +72,7 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'delete' not yet implemented}}
#pragma acc parallel loop auto delete(Var)
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'detach' not yet implemented}}
+ // expected-error at +1{{OpenACC 'detach' clause is not valid on 'parallel loop' directive}}
#pragma acc parallel loop auto detach(Var)
for(unsigned i = 0; i < 5; ++i);
// expected-warning at +1{{OpenACC clause 'device' not yet implemented}}
@@ -189,7 +189,7 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'delete' not yet implemented}}
#pragma acc parallel loop delete(Var) auto
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'detach' not yet implemented}}
+ // expected-error at +1{{OpenACC 'detach' clause is not valid on 'parallel loop' directive}}
#pragma acc parallel loop detach(Var) auto
for(unsigned i = 0; i < 5; ++i);
// expected-warning at +1{{OpenACC clause 'device' not yet implemented}}
@@ -307,7 +307,7 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'delete' not yet implemented}}
#pragma acc parallel loop independent delete(Var)
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'detach' not yet implemented}}
+ // expected-error at +1{{OpenACC 'detach' clause is not valid on 'parallel loop' directive}}
#pragma acc parallel loop independent detach(Var)
for(unsigned i = 0; i < 5; ++i);
// expected-warning at +1{{OpenACC clause 'device' not yet implemented}}
@@ -424,7 +424,7 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'delete' not yet implemented}}
#pragma acc parallel loop delete(Var) independent
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'detach' not yet implemented}}
+ // expected-error at +1{{OpenACC 'detach' clause is not valid on 'parallel loop' directive}}
#pragma acc parallel loop detach(Var) independent
for(unsigned i = 0; i < 5; ++i);
// expected-warning at +1{{OpenACC clause 'device' not yet implemented}}
@@ -550,7 +550,7 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'delete' not yet implemented}}
#pragma acc parallel loop seq delete(Var)
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'detach' not yet implemented}}
+ // expected-error at +1{{OpenACC 'detach' clause is not valid on 'parallel loop' directive}}
#pragma acc parallel loop seq detach(Var)
for(unsigned i = 0; i < 5; ++i);
// expected-warning at +1{{OpenACC clause 'device' not yet implemented}}
@@ -673,7 +673,7 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'delete' not yet implemented}}
#pragma acc parallel loop delete(Var) seq
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'detach' not yet implemented}}
+ // expected-error at +1{{OpenACC 'detach' clause is not valid on 'parallel loop' directive}}
#pragma acc parallel loop detach(Var) seq
for(unsigned i = 0; i < 5; ++i);
// expected-warning at +1{{OpenACC clause 'device' not yet implemented}}
diff --git a/clang/test/SemaOpenACC/combined-construct-device_type-clause.c b/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
index 1a4bd94063e6c5..cc8d8adbdc9f1c 100644
--- a/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
@@ -99,8 +99,7 @@ void uses() {
// expected-note at +1{{previous clause is here}}
#pragma acc serial loop device_type(*) delete(Var)
for(int i = 0; i < 5; ++i);
- // expected-error at +2{{OpenACC clause 'detach' may not follow a 'device_type' clause in a 'kernels loop' construct}}
- // expected-note at +1{{previous clause is here}}
+ // expected-error at +1{{OpenACC 'detach' clause is not valid on 'kernels loop' directive}}
#pragma acc kernels loop device_type(*) detach(Var)
for(int i = 0; i < 5; ++i);
// expected-error at +2{{OpenACC clause 'device' may not follow a 'device_type' clause in a 'parallel loop' construct}}
diff --git a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
index dff8edb7f3dbfd..6c7233e06d7758 100644
--- a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
@@ -103,8 +103,7 @@ void uses() {
// expected-note at +1{{previous clause is here}}
#pragma acc kernels device_type(*) delete(Var)
while(1);
- // expected-error at +2{{OpenACC clause 'detach' may not follow a 'device_type' clause in a 'kernels' construct}}
- // expected-note at +1{{previous clause is here}}
+ // expected-error at +1{{OpenACC 'detach' clause is not valid on 'kernels' directive}}
#pragma acc kernels device_type(*) detach(Var)
while(1);
// expected-error at +2{{OpenACC clause 'device' may not follow a 'device_type' clause in a 'kernels' construct}}
diff --git a/clang/test/SemaOpenACC/data-construct-detach-ast.cpp b/clang/test/SemaOpenACC/data-construct-detach-ast.cpp
new file mode 100644
index 00000000000000..d5096cdc868d08
--- /dev/null
+++ b/clang/test/SemaOpenACC/data-construct-detach-ast.cpp
@@ -0,0 +1,61 @@
+// RUN: %clang_cc1 %s -fopenacc -ast-dump | FileCheck %s
+
+// Test this with PCH.
+// RUN: %clang_cc1 %s -fopenacc -emit-pch -o %t %s
+// RUN: %clang_cc1 %s -fopenacc -include-pch %t -ast-dump-all | FileCheck %s
+
+#ifndef PCH_HELPER
+#define PCH_HELPER
+
+int Global;
+short GlobalArray[5];
+
+void NormalUses(float *PointerParam) {
+ // CHECK: FunctionDecl{{.*}}NormalUses
+ // CHECK: ParmVarDecl
+ // CHECK-NEXT: CompoundStmt
+
+#pragma acc exit data copyout(Global) detach(PointerParam)
+ // CHECK-NEXT: OpenACCExitDataConstruct{{.*}} exit data
+ // CHECK-NEXT: copyout clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Global' 'int'
+ // CHECK-NEXT: detach clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'float *' lvalue ParmVar{{.*}} 'PointerParam' 'float *'
+}
+
+template<typename T>
+void TemplUses(T *t) {
+ // CHECK-NEXT: FunctionTemplateDecl
+ // CHECK-NEXT: TemplateTypeParmDecl{{.*}}typename depth 0 index 0 T
+ // CHECK-NEXT: FunctionDecl{{.*}} TemplUses 'void (T *)'
+ // CHECK-NEXT: ParmVarDecl{{.*}} referenced t 'T *'
+ // CHECK-NEXT: CompoundStmt
+
+#pragma acc exit data copyout(Global) detach(t)
+ // CHECK-NEXT: OpenACCExitDataConstruct{{.*}} exit data
+ // CHECK-NEXT: copyout clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Global' 'int'
+ // CHECK-NEXT: detach clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'T *' lvalue ParmVar{{.*}} 't' 'T *'
+
+
+ // Check the instantiated versions of the above.
+ // CHECK-NEXT: FunctionDecl{{.*}} used TemplUses 'void (int *)' implicit_instantiation
+ // CHECK-NEXT: TemplateArgument type 'int'
+ // CHECK-NEXT: BuiltinType{{.*}} 'int'
+ // CHECK-NEXT: ParmVarDecl{{.*}} used t 'int *'
+ // CHECK-NEXT: CompoundStmt
+
+ // CHECK-NEXT: OpenACCExitDataConstruct{{.*}} exit data
+ // CHECK-NEXT: copyout clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Global' 'int'
+ // CHECK-NEXT: detach clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'int *' lvalue ParmVar{{.*}} 't' 'int *'
+
+}
+
+void Inst() {
+ int i;
+ TemplUses(&i);
+}
+#endif
diff --git a/clang/test/SemaOpenACC/data-construct-detach-clause.c b/clang/test/SemaOpenACC/data-construct-detach-clause.c
new file mode 100644
index 00000000000000..e75c95d99ec078
--- /dev/null
+++ b/clang/test/SemaOpenACC/data-construct-detach-clause.c
@@ -0,0 +1,68 @@
+// 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 'detach' clause, type is 'int'}}
+#pragma acc exit data copyout(LocalInt) detach(LocalInt)
+ ;
+
+ // expected-error at +1{{OpenACC variable is not a valid variable name, sub-array, array element, member of a composite variable, or composite variable member}}
+#pragma acc exit data copyout(LocalInt) detach(&LocalInt)
+ ;
+
+
+ // expected-error at +1{{expected pointer in 'detach' clause, type is 'int[5]'}}
+#pragma acc exit data copyout(LocalInt) detach(Array)
+
+ // expected-error at +1{{expected pointer in 'detach' clause, type is 'int'}}
+#pragma acc exit data copyout(LocalInt) detach(Array[0])
+ ;
+
+ // expected-error at +2{{OpenACC sub-array is not allowed here}}
+ // expected-note at +1{{expected variable of pointer type}}
+#pragma acc exit data copyout(LocalInt) detach(Array[0:1])
+ ;
+
+ // expected-error at +1{{expected pointer in 'detach' clause, type is 'int *[5]'}}
+#pragma acc exit data copyout(LocalInt) detach(PtrArray)
+ ;
+
+#pragma acc exit data copyout(LocalInt) detach(PtrArray[0])
+ ;
+
+ // expected-error at +2{{OpenACC sub-array is not allowed here}}
+ // expected-note at +1{{expected variable of pointer type}}
+#pragma acc exit data copyout(LocalInt) detach(PtrArray[0:1])
+ ;
+
+ // expected-error at +1{{expected pointer in 'detach' clause, type is 'struct S'}}
+#pragma acc exit data copyout(LocalInt) detach(s)
+ ;
+
+ // expected-error at +1{{expected pointer in 'detach' clause, type is 'int'}}
+#pragma acc exit data copyout(LocalInt) detach(s.IntMem)
+ ;
+
+#pragma acc exit data copyout(LocalInt) detach(s.PtrMem)
+ ;
+
+ // expected-error at +1{{OpenACC 'detach' clause is not valid on 'data' directive}}
+#pragma acc data copyin(LocalInt) detach(PtrArray[0])
+ ;
+ // expected-error at +1{{OpenACC 'detach' clause is not valid on 'enter data' directive}}
+#pragma acc enter data copyin(LocalInt) detach(PtrArray[0])
+ // expected-warning at +2{{OpenACC clause 'use_device' not yet implemented}}
+ // expected-error at +1{{OpenACC 'detach' clause is not valid on 'host_data' directive}}
+#pragma acc host_data use_device(LocalInt) detach(PtrArray[0])
+ ;
+}
diff --git a/clang/test/SemaOpenACC/data-construct.cpp b/clang/test/SemaOpenACC/data-construct.cpp
index 0be88c7724e613..1fcf147b0f1b34 100644
--- a/clang/test/SemaOpenACC/data-construct.cpp
+++ b/clang/test/SemaOpenACC/data-construct.cpp
@@ -75,8 +75,7 @@ void AtLeastOneOf() {
#pragma acc exit data copyout(Var)
// expected-warning at +1{{OpenACC clause 'delete' not yet implemented}}
#pragma acc exit data delete(Var)
- // expected-warning at +1{{OpenACC clause 'detach' not yet implemented}}
-#pragma acc exit data detach(Var)
+#pragma acc exit data detach(VarPtr)
// OpenACC TODO: The following 'exit data' directives should diagnose, since
// they don't have at least one of the above clauses.
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 a9ef1a03654a5a..e1e35d6a7b36e0 100644
--- a/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
+++ b/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
@@ -77,7 +77,7 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'delete' not yet implemented}}
#pragma acc loop auto delete(Var)
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'detach' not yet implemented}}
+ // expected-error at +1{{OpenACC 'detach' clause is not valid on 'loop' directive}}
#pragma acc loop auto detach(Var)
for(unsigned i = 0; i < 5; ++i);
// expected-warning at +1{{OpenACC clause 'device' not yet implemented}}
@@ -211,7 +211,7 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'delete' not yet implemented}}
#pragma acc loop delete(Var) auto
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'detach' not yet implemented}}
+ // expected-error at +1{{OpenACC 'detach' clause is not valid on 'loop' directive}}
#pragma acc loop detach(Var) auto
for(unsigned i = 0; i < 5; ++i);
// expected-warning at +1{{OpenACC clause 'device' not yet implemented}}
@@ -346,7 +346,7 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'delete' not yet implemented}}
#pragma acc loop independent delete(Var)
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'detach' not yet implemented}}
+ // expected-error at +1{{OpenACC 'detach' clause is not valid on 'loop' directive}}
#pragma acc loop independent detach(Var)
for(unsigned i = 0; i < 5; ++i);
// expected-warning at +1{{OpenACC clause 'device' not yet implemented}}
@@ -480,7 +480,7 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'delete' not yet implemented}}
#pragma acc loop delete(Var) independent
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'detach' not yet implemented}}
+ // expected-error at +1{{OpenACC 'detach' clause is not valid on 'loop' directive}}
#pragma acc loop detach(Var) independent
for(unsigned i = 0; i < 5; ++i);
// expected-warning at +1{{OpenACC clause 'device' not yet implemented}}
@@ -623,7 +623,7 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'delete' not yet implemented}}
#pragma acc loop seq delete(Var)
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'detach' not yet implemented}}
+ // expected-error at +1{{OpenACC 'detach' clause is not valid on 'loop' directive}}
#pragma acc loop seq detach(Var)
for(unsigned i = 0; i < 5; ++i);
// expected-warning at +1{{OpenACC clause 'device' not yet implemented}}
@@ -763,7 +763,7 @@ void uses() {
// expected-warning at +1{{OpenACC clause 'delete' not yet implemented}}
#pragma acc loop delete(Var) seq
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'detach' not yet implemented}}
+ // expected-error at +1{{OpenACC 'detach' clause is not valid on 'loop' directive}}
#pragma acc loop detach(Var) seq
for(unsigned i = 0; i < 5; ++i);
// expected-warning at +1{{OpenACC clause 'device' not yet implemented}}
diff --git a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
index ad572ad75bca65..94eb2672e8111e 100644
--- a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
@@ -91,8 +91,7 @@ void uses() {
// expected-note at +1{{previous clause is here}}
#pragma acc loop device_type(*) delete(Var)
for(int i = 0; i < 5; ++i);
- // expected-error at +2{{OpenACC clause 'detach' may not follow a 'device_type' clause in a 'loop' construct}}
- // expected-note at +1{{previous clause is here}}
+ // expected-error at +1{{OpenACC 'detach' clause is not valid on 'loop' directive}}
#pragma acc loop device_type(*) detach(Var)
for(int i = 0; i < 5; ++i);
// expected-error at +2{{OpenACC clause 'device' may not follow a 'device_type' clause in a 'loop' construct}}
diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
index c5fdb9065a1c7a..8d6994128f2f07 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -2886,6 +2886,11 @@ void OpenACCClauseEnqueue::VisitCreateClause(const OpenACCCreateClause &C) {
void OpenACCClauseEnqueue::VisitAttachClause(const OpenACCAttachClause &C) {
VisitVarList(C);
}
+
+void OpenACCClauseEnqueue::VisitDetachClause(const OpenACCDetachClause &C) {
+ VisitVarList(C);
+}
+
void OpenACCClauseEnqueue::VisitDevicePtrClause(
const OpenACCDevicePtrClause &C) {
VisitVarList(C);
More information about the cfe-commits
mailing list