[clang] fbb14dd - [OpenACC] Implement 'use_device' clause AST/Sema
via cfe-commits
cfe-commits at lists.llvm.org
Mon Dec 16 09:36:03 PST 2024
Author: erichkeane
Date: 2024-12-16T09:35:57-08:00
New Revision: fbb14dd97702db242a31e1b36ca8a3554a73c212
URL: https://github.com/llvm/llvm-project/commit/fbb14dd97702db242a31e1b36ca8a3554a73c212
DIFF: https://github.com/llvm/llvm-project/commit/fbb14dd97702db242a31e1b36ca8a3554a73c212.diff
LOG: [OpenACC] Implement 'use_device' clause AST/Sema
This is a clause that is only valid on 'host_data' constructs, and
identifies variables which it should use the current device address.
>From a Sema perspective, the only thing novel here is mild changes to
how ActOnVar works for this clause, else this is very much like the rest
of the 'var-list' clauses.
Added:
clang/test/SemaOpenACC/data-construct-use_device-ast.cpp
clang/test/SemaOpenACC/data-construct-use_device-clause.c
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-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-ast.cpp
clang/test/SemaOpenACC/data-construct-async-clause.c
clang/test/SemaOpenACC/data-construct-attach-clause.c
clang/test/SemaOpenACC/data-construct-detach-clause.c
clang/test/SemaOpenACC/data-construct-deviceptr-clause.c
clang/test/SemaOpenACC/data-construct-finalize-clause.c
clang/test/SemaOpenACC/data-construct-if-ast.cpp
clang/test/SemaOpenACC/data-construct-if-clause.c
clang/test/SemaOpenACC/data-construct-if_present-ast.cpp
clang/test/SemaOpenACC/data-construct-if_present-clause.c
clang/test/SemaOpenACC/data-construct-present-clause.c
clang/test/SemaOpenACC/data-construct-wait-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 93053c0e60758e..7a1b17cc4e44e3 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -788,6 +788,27 @@ class OpenACCDeleteClause final
ArrayRef<Expr *> VarList, SourceLocation EndLoc);
};
+class OpenACCUseDeviceClause final
+ : public OpenACCClauseWithVarList,
+ public llvm::TrailingObjects<OpenACCUseDeviceClause, Expr *> {
+
+ OpenACCUseDeviceClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
+ ArrayRef<Expr *> VarList, SourceLocation EndLoc)
+ : OpenACCClauseWithVarList(OpenACCClauseKind::UseDevice, 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::UseDevice;
+ }
+ static OpenACCUseDeviceClause *
+ Create(const ASTContext &C, SourceLocation BeginLoc, SourceLocation LParenLoc,
+ ArrayRef<Expr *> VarList, SourceLocation EndLoc);
+};
class OpenACCNoCreateClause final
: public OpenACCClauseWithVarList,
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 811265151fa0da..77f84a89db2fc9 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12682,6 +12682,9 @@ def err_acc_not_a_var_ref
: Error<"OpenACC variable is not a valid variable name, sub-array, array "
"element,%select{| member of a composite variable,}0 or composite "
"variable member">;
+def err_acc_not_a_var_ref_use_device
+ : Error<"OpenACC variable in 'use_device' clause is not a valid variable "
+ "name or array name">;
def err_acc_typecheck_subarray_value
: Error<"OpenACC sub-array subscripted value is not an array or pointer">;
def err_acc_subarray_function_type
diff --git a/clang/include/clang/Basic/OpenACCClauses.def b/clang/include/clang/Basic/OpenACCClauses.def
index 600510e6980dae..c7ffac391e2026 100644
--- a/clang/include/clang/Basic/OpenACCClauses.def
+++ b/clang/include/clang/Basic/OpenACCClauses.def
@@ -58,6 +58,7 @@ VISIT_CLAUSE(Reduction)
VISIT_CLAUSE(Self)
VISIT_CLAUSE(Seq)
VISIT_CLAUSE(Tile)
+VISIT_CLAUSE(UseDevice)
VISIT_CLAUSE(Vector)
VISIT_CLAUSE(VectorLength)
VISIT_CLAUSE(Wait)
diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h
index 58137d3a7e3f73..78056d03680017 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -400,6 +400,7 @@ class SemaOpenACC : public SemaBase {
ClauseKind == OpenACCClauseKind::PresentOrCreate ||
ClauseKind == OpenACCClauseKind::Attach ||
ClauseKind == OpenACCClauseKind::Delete ||
+ ClauseKind == OpenACCClauseKind::UseDevice ||
ClauseKind == OpenACCClauseKind::Detach ||
ClauseKind == OpenACCClauseKind::DevicePtr ||
ClauseKind == OpenACCClauseKind::Reduction ||
@@ -538,6 +539,7 @@ class SemaOpenACC : public SemaBase {
ClauseKind == OpenACCClauseKind::PresentOrCreate ||
ClauseKind == OpenACCClauseKind::Attach ||
ClauseKind == OpenACCClauseKind::Delete ||
+ ClauseKind == OpenACCClauseKind::UseDevice ||
ClauseKind == OpenACCClauseKind::Detach ||
ClauseKind == OpenACCClauseKind::DevicePtr ||
ClauseKind == OpenACCClauseKind::FirstPrivate) &&
@@ -576,6 +578,7 @@ class SemaOpenACC : public SemaBase {
ClauseKind == OpenACCClauseKind::PresentOrCreate ||
ClauseKind == OpenACCClauseKind::Attach ||
ClauseKind == OpenACCClauseKind::Delete ||
+ ClauseKind == OpenACCClauseKind::UseDevice ||
ClauseKind == OpenACCClauseKind::Detach ||
ClauseKind == OpenACCClauseKind::DevicePtr ||
ClauseKind == OpenACCClauseKind::FirstPrivate) &&
diff --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp
index f836d30561e33b..fbc9f6d15fa7bf 100644
--- a/clang/lib/AST/OpenACCClause.cpp
+++ b/clang/lib/AST/OpenACCClause.cpp
@@ -33,6 +33,7 @@ bool OpenACCClauseWithVarList::classof(const OpenACCClause *C) {
OpenACCFirstPrivateClause::classof(C) ||
OpenACCDevicePtrClause::classof(C) ||
OpenACCDeleteClause::classof(C) ||
+ OpenACCUseDeviceClause::classof(C) ||
OpenACCDetachClause::classof(C) || OpenACCAttachClause::classof(C) ||
OpenACCNoCreateClause::classof(C) ||
OpenACCPresentClause::classof(C) || OpenACCCopyClause::classof(C) ||
@@ -298,6 +299,16 @@ OpenACCDeleteClause *OpenACCDeleteClause::Create(const ASTContext &C,
return new (Mem) OpenACCDeleteClause(BeginLoc, LParenLoc, VarList, EndLoc);
}
+OpenACCUseDeviceClause *OpenACCUseDeviceClause::Create(const ASTContext &C,
+ SourceLocation BeginLoc,
+ SourceLocation LParenLoc,
+ ArrayRef<Expr *> VarList,
+ SourceLocation EndLoc) {
+ void *Mem = C.Allocate(
+ OpenACCUseDeviceClause::totalSizeToAlloc<Expr *>(VarList.size()));
+ return new (Mem) OpenACCUseDeviceClause(BeginLoc, LParenLoc, VarList, EndLoc);
+}
+
OpenACCDevicePtrClause *OpenACCDevicePtrClause::Create(const ASTContext &C,
SourceLocation BeginLoc,
SourceLocation LParenLoc,
@@ -581,6 +592,14 @@ void OpenACCClausePrinter::VisitDeleteClause(const OpenACCDeleteClause &C) {
OS << ")";
}
+void OpenACCClausePrinter::VisitUseDeviceClause(
+ const OpenACCUseDeviceClause &C) {
+ OS << "use_device(";
+ 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 6160a69832e8aa..1fb238720ffb13 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -2515,6 +2515,11 @@ class OpenACCClauseProfiler
}
}
+ void VisitClauseWithVarList(const OpenACCClauseWithVarList &Clause) {
+ for (auto *E : Clause.getVarList())
+ Profiler.VisitStmt(E);
+ }
+
#define VISIT_CLAUSE(CLAUSE_NAME) \
void Visit##CLAUSE_NAME##Clause(const OpenACC##CLAUSE_NAME##Clause &Clause);
@@ -2532,25 +2537,21 @@ void OpenACCClauseProfiler::VisitIfClause(const OpenACCIfClause &Clause) {
}
void OpenACCClauseProfiler::VisitCopyClause(const OpenACCCopyClause &Clause) {
- for (auto *E : Clause.getVarList())
- Profiler.VisitStmt(E);
+ VisitClauseWithVarList(Clause);
}
void OpenACCClauseProfiler::VisitCopyInClause(
const OpenACCCopyInClause &Clause) {
- for (auto *E : Clause.getVarList())
- Profiler.VisitStmt(E);
+ VisitClauseWithVarList(Clause);
}
void OpenACCClauseProfiler::VisitCopyOutClause(
const OpenACCCopyOutClause &Clause) {
- for (auto *E : Clause.getVarList())
- Profiler.VisitStmt(E);
+ VisitClauseWithVarList(Clause);
}
void OpenACCClauseProfiler::VisitCreateClause(
const OpenACCCreateClause &Clause) {
- for (auto *E : Clause.getVarList())
- Profiler.VisitStmt(E);
+ VisitClauseWithVarList(Clause);
}
void OpenACCClauseProfiler::VisitSelfClause(const OpenACCSelfClause &Clause) {
@@ -2589,50 +2590,47 @@ void OpenACCClauseProfiler::VisitCollapseClause(
void OpenACCClauseProfiler::VisitPrivateClause(
const OpenACCPrivateClause &Clause) {
- for (auto *E : Clause.getVarList())
- Profiler.VisitStmt(E);
+ VisitClauseWithVarList(Clause);
}
void OpenACCClauseProfiler::VisitFirstPrivateClause(
const OpenACCFirstPrivateClause &Clause) {
- for (auto *E : Clause.getVarList())
- Profiler.VisitStmt(E);
+ VisitClauseWithVarList(Clause);
}
void OpenACCClauseProfiler::VisitAttachClause(
const OpenACCAttachClause &Clause) {
- for (auto *E : Clause.getVarList())
- Profiler.VisitStmt(E);
+ VisitClauseWithVarList(Clause);
}
void OpenACCClauseProfiler::VisitDetachClause(
const OpenACCDetachClause &Clause) {
- for (auto *E : Clause.getVarList())
- Profiler.VisitStmt(E);
+ VisitClauseWithVarList(Clause);
}
void OpenACCClauseProfiler::VisitDeleteClause(
const OpenACCDeleteClause &Clause) {
- for (auto *E : Clause.getVarList())
- Profiler.VisitStmt(E);
+ VisitClauseWithVarList(Clause);
}
void OpenACCClauseProfiler::VisitDevicePtrClause(
const OpenACCDevicePtrClause &Clause) {
- for (auto *E : Clause.getVarList())
- Profiler.VisitStmt(E);
+ VisitClauseWithVarList(Clause);
}
void OpenACCClauseProfiler::VisitNoCreateClause(
const OpenACCNoCreateClause &Clause) {
- for (auto *E : Clause.getVarList())
- Profiler.VisitStmt(E);
+ VisitClauseWithVarList(Clause);
}
void OpenACCClauseProfiler::VisitPresentClause(
const OpenACCPresentClause &Clause) {
- for (auto *E : Clause.getVarList())
- Profiler.VisitStmt(E);
+ VisitClauseWithVarList(Clause);
+}
+
+void OpenACCClauseProfiler::VisitUseDeviceClause(
+ const OpenACCUseDeviceClause &Clause) {
+ VisitClauseWithVarList(Clause);
}
void OpenACCClauseProfiler::VisitVectorLengthClause(
@@ -2684,8 +2682,7 @@ void OpenACCClauseProfiler::VisitGangClause(const OpenACCGangClause &Clause) {
void OpenACCClauseProfiler::VisitReductionClause(
const OpenACCReductionClause &Clause) {
- for (auto *E : Clause.getVarList())
- Profiler.VisitStmt(E);
+ VisitClauseWithVarList(Clause);
}
} // namespace
diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp
index 6040f34a4b9a5f..b5af10dd00b77c 100644
--- a/clang/lib/AST/TextNodeDumper.cpp
+++ b/clang/lib/AST/TextNodeDumper.cpp
@@ -425,6 +425,7 @@ void TextNodeDumper::Visit(const OpenACCClause *C) {
case OpenACCClauseKind::Seq:
case OpenACCClauseKind::Tile:
case OpenACCClauseKind::Worker:
+ case OpenACCClauseKind::UseDevice:
case OpenACCClauseKind::Vector:
case OpenACCClauseKind::VectorLength:
// The condition expression will be printed as a part of the 'children',
diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp
index 5130159f5d8ac8..570dba811aca86 100644
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@@ -1002,13 +1002,13 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
case OpenACCClauseKind::DeviceResident:
case OpenACCClauseKind::Host:
case OpenACCClauseKind::Link:
- case OpenACCClauseKind::UseDevice:
ParseOpenACCVarList(ClauseKind);
break;
case OpenACCClauseKind::Attach:
case OpenACCClauseKind::Delete:
case OpenACCClauseKind::Detach:
case OpenACCClauseKind::DevicePtr:
+ case OpenACCClauseKind::UseDevice:
ParsedClause.setVarListDetails(ParseOpenACCVarList(ClauseKind),
/*IsReadOnly=*/false, /*IsZero=*/false);
break;
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 7af209aa155968..79822bf583195e 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -442,6 +442,15 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
return false;
}
}
+
+ case OpenACCClauseKind::UseDevice: {
+ switch (DirectiveKind) {
+ case OpenACCDirectiveKind::HostData:
+ return true;
+ default:
+ return false;
+ }
+ }
}
default:
@@ -1085,6 +1094,14 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitDeleteClause(
Clause.getEndLoc());
}
+OpenACCClause *SemaOpenACCClauseVisitor::VisitUseDeviceClause(
+ SemaOpenACC::OpenACCParsedClause &Clause) {
+ // ActOnVar ensured that everything is a valid variable or array, so nothing
+ // left to do here.
+ return OpenACCUseDeviceClause::Create(
+ Ctx, Clause.getBeginLoc(), Clause.getLParenLoc(), Clause.getVarList(),
+ Clause.getEndLoc());
+}
OpenACCClause *SemaOpenACCClauseVisitor::VisitDevicePtrClause(
SemaOpenACC::OpenACCParsedClause &Clause) {
@@ -2431,6 +2448,15 @@ bool SemaOpenACC::CheckVarIsPointerType(OpenACCClauseKind ClauseKind,
ExprResult SemaOpenACC::ActOnVar(OpenACCClauseKind CK, Expr *VarExpr) {
Expr *CurVarExpr = VarExpr->IgnoreParenImpCasts();
+ // 'use_device' doesn't allow array subscript or array sections.
+ // OpenACC3.3 2.8:
+ // A 'var' in a 'use_device' clause must be the name of a variable or array.
+ if (CK == OpenACCClauseKind::UseDevice &&
+ isa<ArraySectionExpr, ArraySubscriptExpr>(CurVarExpr)) {
+ Diag(VarExpr->getExprLoc(), diag::err_acc_not_a_var_ref_use_device);
+ return ExprError();
+ }
+
// Sub-arrays/subscript-exprs are fine as long as the base is a
// VarExpr/MemberExpr. So strip all of those off.
while (isa<ArraySectionExpr, ArraySubscriptExpr>(CurVarExpr)) {
@@ -2451,16 +2477,20 @@ ExprResult SemaOpenACC::ActOnVar(OpenACCClauseKind CK, Expr *VarExpr) {
// If CK is a Reduction, this special cases for OpenACC3.3 2.5.15: "A var in a
// reduction clause must be a scalar variable name, an aggregate variable
// name, an array element, or a subarray.
- // A MemberExpr that references a Field is valid.
- if (CK != OpenACCClauseKind::Reduction) {
+ // If CK is a 'use_device', this also isn't valid, as it isn' the name of a
+ // variable or array.
+ // A MemberExpr that references a Field is valid for other clauses.
+ if (CK != OpenACCClauseKind::Reduction &&
+ CK != OpenACCClauseKind::UseDevice) {
if (const auto *ME = dyn_cast<MemberExpr>(CurVarExpr)) {
if (isa<FieldDecl>(ME->getMemberDecl()->getCanonicalDecl()))
return VarExpr;
}
}
- // Referring to 'this' is always OK.
- if (isa<CXXThisExpr>(CurVarExpr))
+ // Referring to 'this' is ok for the most part, but for 'use_device' doesn't
+ // fall into 'variable or array name'
+ if (CK != OpenACCClauseKind::UseDevice && isa<CXXThisExpr>(CurVarExpr))
return VarExpr;
// Nothing really we can do here, as these are dependent. So just return they
@@ -2475,8 +2505,11 @@ ExprResult SemaOpenACC::ActOnVar(OpenACCClauseKind CK, Expr *VarExpr) {
if (isa<RecoveryExpr>(CurVarExpr))
return ExprError();
- Diag(VarExpr->getExprLoc(), diag::err_acc_not_a_var_ref)
- << (CK != OpenACCClauseKind::Reduction);
+ if (CK == OpenACCClauseKind::UseDevice)
+ Diag(VarExpr->getExprLoc(), diag::err_acc_not_a_var_ref_use_device);
+ else
+ Diag(VarExpr->getExprLoc(), diag::err_acc_not_a_var_ref)
+ << (CK != OpenACCClauseKind::Reduction);
return ExprError();
}
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index c33648ca0e34b8..c9bb079a9bcb34 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -11788,6 +11788,17 @@ void OpenACCClauseTransform<Derived>::VisitDeleteClause(
ParsedClause.getEndLoc());
}
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitUseDeviceClause(
+ const OpenACCUseDeviceClause &C) {
+ ParsedClause.setVarListDetails(VisitVarList(C.getVarList()),
+ /*IsReadOnly=*/false, /*IsZero=*/false);
+ NewClause = OpenACCUseDeviceClause::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 741bae684cffe3..21f6b2ecc58c4f 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -12441,6 +12441,12 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
return OpenACCDeleteClause::Create(getContext(), BeginLoc, LParenLoc,
VarList, EndLoc);
}
+ case OpenACCClauseKind::UseDevice: {
+ SourceLocation LParenLoc = readSourceLocation();
+ llvm::SmallVector<Expr *> VarList = readOpenACCVarList();
+ return OpenACCUseDeviceClause::Create(getContext(), BeginLoc, LParenLoc,
+ VarList, EndLoc);
+ }
case OpenACCClauseKind::DevicePtr: {
SourceLocation LParenLoc = readSourceLocation();
llvm::SmallVector<Expr *> VarList = readOpenACCVarList();
@@ -12583,7 +12589,6 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
}
case OpenACCClauseKind::NoHost:
- case OpenACCClauseKind::UseDevice:
case OpenACCClauseKind::Device:
case OpenACCClauseKind::DeviceResident:
case OpenACCClauseKind::Host:
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index 9517bad4070dfb..6db2262a7952ec 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -8368,6 +8368,12 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
writeOpenACCVarList(DC);
return;
}
+ case OpenACCClauseKind::UseDevice: {
+ const auto *UDC = cast<OpenACCUseDeviceClause>(C);
+ writeSourceLocation(UDC->getLParenLoc());
+ writeOpenACCVarList(UDC);
+ return;
+ }
case OpenACCClauseKind::DevicePtr: {
const auto *DPC = cast<OpenACCDevicePtrClause>(C);
writeSourceLocation(DPC->getLParenLoc());
@@ -8511,7 +8517,6 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
}
case OpenACCClauseKind::NoHost:
- case OpenACCClauseKind::UseDevice:
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 f03f96239ab4c6..7ad0a43e322a21 100644
--- a/clang/test/AST/ast-print-openacc-data-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-data-construct.cpp
@@ -20,8 +20,7 @@ void foo() {
// CHECK: #pragma acc exit data copyout(Var)
#pragma acc exit data copyout(Var)
;
-// CHECK: #pragma acc host_data
-// CHECK-NOT: use_device(Var)
+// CHECK: #pragma acc host_data use_device(Var)
#pragma acc host_data use_device(Var)
;
@@ -38,7 +37,7 @@ void foo() {
// CHECK: #pragma acc exit data copyout(Var) if(i == array[1])
#pragma acc exit data copyout(Var) if(i == array[1])
;
-// CHECK: #pragma acc host_data if(i == array[1])
+// CHECK: #pragma acc host_data use_device(Var) if(i == array[1])
#pragma acc host_data use_device(Var) if(i == array[1])
;
@@ -114,7 +113,7 @@ void foo() {
// CHECK: #pragma acc exit data copyout(i) finalize
#pragma acc exit data copyout(i) finalize
-// CHECK: #pragma acc host_data if_present
+// CHECK: #pragma acc host_data use_device(i) if_present
#pragma acc host_data use_device(i) if_present
;
// CHECK: #pragma acc exit data copyout(i) detach(iPtr, arrayPtr[0])
@@ -126,4 +125,8 @@ void foo() {
// CHECK: #pragma acc exit data copyout(i) delete(i, array[1], array, array[1:2])
#pragma acc exit data copyout(i) delete(i, array[1], array, array[1:2])
+
+// CHECK: #pragma acc host_data use_device(i)
+#pragma acc host_data use_device(i)
+ ;
}
diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c
index e583fb3897998d..487dc79f538085 100644
--- a/clang/test/ParserOpenACC/parse-clauses.c
+++ b/clang/test/ParserOpenACC/parse-clauses.c
@@ -443,13 +443,16 @@ void VarListClauses() {
#pragma acc serial present_or_copy(HasMem.MemArr[3:])
for(int i = 0; i < 5;++i) {}
- // expected-error at +2{{expected ','}}
- // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented, clause ignored}}
-#pragma acc serial use_device(s.array[s.value] s.array[s.value :5] ), self
+ // expected-error at +2 2{{OpenACC variable in 'use_device' clause is not a valid variable name or array name}}
+ // expected-error at +1{{expected ','}}
+#pragma acc host_data use_device(s.array[s.value] s.array[s.value :5] ), if_present
for(int i = 0; i < 5;++i) {}
- // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented, clause ignored}}
-#pragma acc serial use_device(s.array[s.value : 5]), self
+ // expected-error at +1{{OpenACC variable in 'use_device' clause is not a valid variable name or array name}}
+#pragma acc host_data use_device(s.array[s.value : 5]), if_present
+ for(int i = 0; i < 5;++i) {}
+
+#pragma acc host_data use_device(HasMem), if_present
for(int i = 0; i < 5;++i) {}
// expected-error at +1{{expected ','}}
@@ -517,15 +520,6 @@ void VarListClauses() {
#pragma acc exit data delete(s.array[s.value : 5], s.value),async
for(int i = 0; i < 5;++i) {}
- // expected-error at +2{{expected ','}}
- // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented, clause ignored}}
-#pragma acc exit data use_device(s.array[s.value] s.array[s.value :5] ),async
- for(int i = 0; i < 5;++i) {}
-
- // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented, clause ignored}}
-#pragma acc exit data use_device(s.array[s.value : 5], s.value), async
- for(int i = 0; i < 5;++i) {}
-
// expected-error at +2{{expected ','}}
// expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented, clause ignored}}
#pragma acc serial device_resident(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 a9f6f1e6b9e3a3..a71d25ce61bed4 100644
--- a/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
+++ b/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
@@ -64,7 +64,7 @@ void uses() {
// expected-warning at +1{{OpenACC clause name 'present_or_copy' is a deprecated clause name and is now an alias for 'copy'}}
#pragma acc parallel loop auto present_or_copy(Var)
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
+ // expected-error at +1{{OpenACC 'use_device' clause is not valid on 'parallel loop' directive}}
#pragma acc parallel loop auto use_device(Var)
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop auto attach(VarPtr)
@@ -181,7 +181,7 @@ void uses() {
// expected-warning at +1{{OpenACC clause name 'present_or_copy' is a deprecated clause name and is now an alias for 'copy'}}
#pragma acc parallel loop present_or_copy(Var) auto
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
+ // expected-error at +1{{OpenACC 'use_device' clause is not valid on 'parallel loop' directive}}
#pragma acc parallel loop use_device(Var) auto
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop attach(VarPtr) auto
@@ -299,7 +299,7 @@ void uses() {
// expected-warning at +1{{OpenACC clause name 'present_or_copy' is a deprecated clause name and is now an alias for 'copy'}}
#pragma acc parallel loop independent present_or_copy(Var)
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
+ // expected-error at +1{{OpenACC 'use_device' clause is not valid on 'parallel loop' directive}}
#pragma acc parallel loop independent use_device(Var)
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop independent attach(VarPtr)
@@ -416,7 +416,7 @@ void uses() {
// expected-warning at +1{{OpenACC clause name 'present_or_copy' is a deprecated clause name and is now an alias for 'copy'}}
#pragma acc parallel loop present_or_copy(Var) independent
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
+ // expected-error at +1{{OpenACC 'use_device' clause is not valid on 'parallel loop' directive}}
#pragma acc parallel loop use_device(Var) independent
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop attach(VarPtr) independent
@@ -542,7 +542,7 @@ void uses() {
// expected-warning at +1{{OpenACC clause name 'present_or_copy' is a deprecated clause name and is now an alias for 'copy'}}
#pragma acc parallel loop seq present_or_copy(Var)
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
+ // expected-error at +1{{OpenACC 'use_device' clause is not valid on 'parallel loop' directive}}
#pragma acc parallel loop seq use_device(Var)
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop seq attach(VarPtr)
@@ -665,7 +665,7 @@ void uses() {
// expected-warning at +1{{OpenACC clause name 'present_or_copy' is a deprecated clause name and is now an alias for 'copy'}}
#pragma acc parallel loop present_or_copy(Var) seq
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
+ // expected-error at +1{{OpenACC 'use_device' clause is not valid on 'parallel loop' directive}}
#pragma acc parallel loop use_device(Var) seq
for(unsigned i = 0; i < 5; ++i);
#pragma acc parallel loop attach(VarPtr) seq
diff --git a/clang/test/SemaOpenACC/combined-construct-device_type-clause.c b/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
index 4526a11eeb9079..40339941f51a9c 100644
--- a/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
@@ -87,8 +87,7 @@ void uses() {
// expected-note at +1{{previous clause is here}}
#pragma acc serial loop device_type(*) present_or_copy(Var)
for(int i = 0; i < 5; ++i);
- // expected-error at +2{{OpenACC clause 'use_device' 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 'use_device' clause is not valid on 'kernels loop' directive}}
#pragma acc kernels loop device_type(*) use_device(Var)
for(int i = 0; i < 5; ++i);
// expected-error at +2{{OpenACC clause 'attach' 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 6f46e615f43c9d..cccb50ee55582d 100644
--- a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
@@ -91,8 +91,7 @@ void uses() {
// expected-note at +1{{previous clause is here}}
#pragma acc kernels device_type(*) present_or_copy(Var)
while(1);
- // expected-error at +2{{OpenACC clause 'use_device' may not follow a 'device_type' clause in a 'kernels' construct}}
- // expected-note at +1{{previous clause is here}}
+ // expected-error at +1{{OpenACC 'use_device' clause is not valid on 'kernels' directive}}
#pragma acc kernels device_type(*) use_device(Var)
while(1);
// expected-error at +2{{OpenACC clause 'attach' may not follow a 'device_type' clause in a 'kernels' construct}}
diff --git a/clang/test/SemaOpenACC/data-construct-ast.cpp b/clang/test/SemaOpenACC/data-construct-ast.cpp
index f299fd04581a72..5abd142b237a21 100644
--- a/clang/test/SemaOpenACC/data-construct-ast.cpp
+++ b/clang/test/SemaOpenACC/data-construct-ast.cpp
@@ -36,6 +36,8 @@ void NormalFunc() {
#pragma acc host_data use_device(Var)
while (Var);
// CHECK-NEXT: OpenACCHostDataConstruct{{.*}} host_data
+ // CHECK-NEXT: use_device clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Var' 'int'
// CHECK-NEXT: WhileStmt
// CHECK: NullStmt
}
@@ -68,6 +70,8 @@ void TemplFunc() {
#pragma acc host_data use_device(Var)
while (Var);
// CHECK-NEXT: OpenACCHostDataConstruct{{.*}} host_data
+ // CHECK-NEXT: use_device clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Var' 'T'
// CHECK-NEXT: WhileStmt
// CHECK: NullStmt
@@ -94,6 +98,8 @@ void TemplFunc() {
// CHECK-NEXT: DeclRefExpr{{.*}}'Var' 'int'
// CHECK-NEXT: OpenACCHostDataConstruct{{.*}} host_data
+ // CHECK-NEXT: use_device clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Var' 'int'
// CHECK-NEXT: WhileStmt
// CHECK: NullStmt
}
diff --git a/clang/test/SemaOpenACC/data-construct-async-clause.c b/clang/test/SemaOpenACC/data-construct-async-clause.c
index 7173b2f0be7dd9..3c9fbae0d9875d 100644
--- a/clang/test/SemaOpenACC/data-construct-async-clause.c
+++ b/clang/test/SemaOpenACC/data-construct-async-clause.c
@@ -9,7 +9,6 @@ void Test() {
;
#pragma acc enter data copyin(I) async(I)
#pragma acc exit data copyout(I) async(I)
- // expected-warning at +2{{OpenACC clause 'use_device' not yet implemented}}
// expected-error at +1{{OpenACC 'async' clause is not valid on 'host_data' directive}}
#pragma acc host_data use_device(I) async(I)
;
@@ -21,7 +20,6 @@ void Test() {
#pragma acc enter data copyin(NC) async(NC)
// expected-error at +1{{OpenACC clause 'async' requires expression of integer type ('struct NotConvertible' invalid)}}
#pragma acc exit data copyout(NC) async(NC)
- // expected-warning at +2{{OpenACC clause 'use_device' not yet implemented}}
// expected-error at +1{{OpenACC clause 'async' requires expression of integer type ('struct NotConvertible' invalid)}}
#pragma acc host_data use_device(NC) async(NC)
;
diff --git a/clang/test/SemaOpenACC/data-construct-attach-clause.c b/clang/test/SemaOpenACC/data-construct-attach-clause.c
index 49a708e49d24b2..0bc02563ff6953 100644
--- a/clang/test/SemaOpenACC/data-construct-attach-clause.c
+++ b/clang/test/SemaOpenACC/data-construct-attach-clause.c
@@ -58,7 +58,6 @@ void uses() {
// expected-error at +1{{OpenACC 'attach' clause is not valid on 'exit data' directive}}
#pragma acc exit data copyout(LocalInt) attach(PtrArray[0])
- // expected-warning at +2{{OpenACC clause 'use_device' not yet implemented}}
// expected-error at +1{{OpenACC 'attach' clause is not valid on 'host_data' directive}}
#pragma acc host_data use_device(LocalInt) attach(PtrArray[0])
;
diff --git a/clang/test/SemaOpenACC/data-construct-detach-clause.c b/clang/test/SemaOpenACC/data-construct-detach-clause.c
index e75c95d99ec078..edcc80ac362362 100644
--- a/clang/test/SemaOpenACC/data-construct-detach-clause.c
+++ b/clang/test/SemaOpenACC/data-construct-detach-clause.c
@@ -61,7 +61,6 @@ void uses() {
;
// 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-deviceptr-clause.c b/clang/test/SemaOpenACC/data-construct-deviceptr-clause.c
index d7869e965c5b6a..70ccd3b7aec953 100644
--- a/clang/test/SemaOpenACC/data-construct-deviceptr-clause.c
+++ b/clang/test/SemaOpenACC/data-construct-deviceptr-clause.c
@@ -61,7 +61,6 @@ void uses() {
#pragma acc enter data copyin(LocalInt) deviceptr(LocalInt)
// expected-error at +1{{OpenACC 'deviceptr' clause is not valid on 'exit data' directive}}
#pragma acc exit data copyout(LocalInt) deviceptr(LocalInt)
- // expected-warning at +2{{OpenACC clause 'use_device' not yet implemented}}
// expected-error at +1{{OpenACC 'deviceptr' clause is not valid on 'host_data' directive}}
#pragma acc host_data use_device(LocalInt) deviceptr(LocalInt)
;
diff --git a/clang/test/SemaOpenACC/data-construct-finalize-clause.c b/clang/test/SemaOpenACC/data-construct-finalize-clause.c
index b2b4ada0e42ed9..252b26708cd811 100644
--- a/clang/test/SemaOpenACC/data-construct-finalize-clause.c
+++ b/clang/test/SemaOpenACC/data-construct-finalize-clause.c
@@ -13,7 +13,6 @@ void Test() {
// finalize is valid only on exit data, otherwise has no other rules.
#pragma acc exit data copyout(I) finalize
;
- // expected-warning at +2{{OpenACC clause 'use_device' not yet implemented}}
// expected-error at +1{{OpenACC 'finalize' clause is not valid on 'host_data' directive}}
#pragma acc host_data use_device(I) finalize
;
diff --git a/clang/test/SemaOpenACC/data-construct-if-ast.cpp b/clang/test/SemaOpenACC/data-construct-if-ast.cpp
index 3b810a724c51ee..9ceee4e1c07497 100644
--- a/clang/test/SemaOpenACC/data-construct-if-ast.cpp
+++ b/clang/test/SemaOpenACC/data-construct-if-ast.cpp
@@ -72,6 +72,8 @@ void TemplFunc() {
#pragma acc host_data use_device(Global) if(T::BC)
;
// CHECK-NEXT: OpenACCHostDataConstruct{{.*}}host_data
+ // CHECK-NEXT: use_device clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Global' 'int'
// CHECK-NEXT: if clause
// CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue
// CHECK-NEXT: NestedNameSpecifier TypeSpec 'T'
@@ -116,6 +118,8 @@ void TemplFunc() {
// CHECK-NEXT: NullStmt
// CHECK-NEXT: OpenACCHostDataConstruct{{.*}}host_data
+ // CHECK-NEXT: use_device clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'Global' 'int'
// CHECK-NEXT: if clause
// CHECK-NEXT: ImplicitCastExpr{{.*}} 'bool' <UserDefinedConversion>
// CHECK-NEXT: CXXMemberCallExpr{{.*}} 'bool'
diff --git a/clang/test/SemaOpenACC/data-construct-if-clause.c b/clang/test/SemaOpenACC/data-construct-if-clause.c
index 0a7989e80fc373..f22452d2c34a41 100644
--- a/clang/test/SemaOpenACC/data-construct-if-clause.c
+++ b/clang/test/SemaOpenACC/data-construct-if-clause.c
@@ -20,10 +20,8 @@ void Foo() {
// expected-note at +1{{previous clause is here}}
#pragma acc exit data copyout(Var) if(1) if (2)
- // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
#pragma acc host_data use_device(Var) if(1)
;
- // expected-warning at +3{{OpenACC clause 'use_device' not yet implemented}}
// expected-error at +2{{OpenACC 'if' clause cannot appear more than once on a 'host_data' directive}}
// expected-note at +1{{previous clause is here}}
#pragma acc host_data use_device(Var) if(1) if (2)
diff --git a/clang/test/SemaOpenACC/data-construct-if_present-ast.cpp b/clang/test/SemaOpenACC/data-construct-if_present-ast.cpp
index d2ef6822819655..8dab6bcce58d9c 100644
--- a/clang/test/SemaOpenACC/data-construct-if_present-ast.cpp
+++ b/clang/test/SemaOpenACC/data-construct-if_present-ast.cpp
@@ -17,6 +17,8 @@ void Uses() {
#pragma acc host_data use_device(I) if_present
;
// CHECK-NEXT: OpenACCHostDataConstruct{{.*}}host_data
+ // CHECK-NEXT: use_device clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'I' 'int'
// CHECK-NEXT: if_present clause
// CHECK-NEXT: NullStmt
}
@@ -35,6 +37,8 @@ void TemplUses() {
#pragma acc host_data use_device(I) if_present
;
// CHECK-NEXT: OpenACCHostDataConstruct{{.*}}host_data
+ // CHECK-NEXT: use_device clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'I' 'T'
// CHECK-NEXT: if_present clause
// CHECK-NEXT: NullStmt
@@ -48,6 +52,8 @@ void TemplUses() {
// CHECK-NEXT: VarDecl
// CHECK-NEXT: OpenACCHostDataConstruct{{.*}}host_data
+ // CHECK-NEXT: use_device clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'I' 'int'
// CHECK-NEXT: if_present clause
// CHECK-NEXT: NullStmt
}
diff --git a/clang/test/SemaOpenACC/data-construct-if_present-clause.c b/clang/test/SemaOpenACC/data-construct-if_present-clause.c
index ce92ec024b4f99..b1290cdccca5f1 100644
--- a/clang/test/SemaOpenACC/data-construct-if_present-clause.c
+++ b/clang/test/SemaOpenACC/data-construct-if_present-clause.c
@@ -13,7 +13,6 @@ void Test() {
// expected-error at +1{{OpenACC 'if_present' clause is not valid on 'exit data' directive}}
#pragma acc exit data copyout(I) if_present
;
- // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
#pragma acc host_data use_device(I) if_present
;
}
diff --git a/clang/test/SemaOpenACC/data-construct-present-clause.c b/clang/test/SemaOpenACC/data-construct-present-clause.c
index 3128f532faf3f0..b889230d177cdc 100644
--- a/clang/test/SemaOpenACC/data-construct-present-clause.c
+++ b/clang/test/SemaOpenACC/data-construct-present-clause.c
@@ -52,7 +52,6 @@ void uses(int IntParam, short *PointerParam, float ArrayParam[5], Complete Compo
#pragma acc enter data copyin(LocalInt) present(LocalInt)
// expected-error at +1{{OpenACC 'present' clause is not valid on 'exit data' directive}}
#pragma acc exit data copyout(LocalInt) present(LocalInt)
- // expected-warning at +2{{OpenACC clause 'use_device' not yet implemented}}
// expected-error at +1{{OpenACC 'present' clause is not valid on 'host_data' directive}}
#pragma acc host_data use_device(LocalInt) present(LocalInt)
;
diff --git a/clang/test/SemaOpenACC/data-construct-use_device-ast.cpp b/clang/test/SemaOpenACC/data-construct-use_device-ast.cpp
new file mode 100644
index 00000000000000..b8cc30e14671fa
--- /dev/null
+++ b/clang/test/SemaOpenACC/data-construct-use_device-ast.cpp
@@ -0,0 +1,58 @@
+// 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 host_data use_device(GlobalArray, PointerParam)
+ ;
+ // CHECK-NEXT: OpenACCHostDataConstruct{{.*}} host_data
+ // CHECK-NEXT: use_device clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'short[5]' lvalue Var{{.*}}'GlobalArray' 'short[5]'
+ // CHECK-NEXT: DeclRefExpr{{.*}}'float *' lvalue ParmVar{{.*}}'PointerParam' 'float *'
+ // CHECK-NEXT: NullStmt
+}
+
+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 host_data use_device(t)
+ ;
+ // CHECK-NEXT: OpenACCHostDataConstruct{{.*}} host_data
+ // CHECK-NEXT: use_device clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'T' lvalue ParmVar{{.*}} 't' 'T'
+ // CHECK-NEXT: NullStmt
+
+ // 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: OpenACCHostDataConstruct{{.*}} host_data
+ // CHECK-NEXT: use_device clause
+ // CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue ParmVar{{.*}} 't' 'int'
+ // CHECK-NEXT: NullStmt
+}
+
+void Inst() {
+ int i;
+ TemplUses(i);
+}
+#endif
diff --git a/clang/test/SemaOpenACC/data-construct-use_device-clause.c b/clang/test/SemaOpenACC/data-construct-use_device-clause.c
new file mode 100644
index 00000000000000..d0f74585759cff
--- /dev/null
+++ b/clang/test/SemaOpenACC/data-construct-use_device-clause.c
@@ -0,0 +1,62 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+typedef struct IsComplete {
+ struct S { int A; } CompositeMember;
+ int ScalarMember;
+ float ArrayMember[5];
+ void *PointerMember;
+} Complete;
+void uses(int IntParam, short *PointerParam, float ArrayParam[5], Complete CompositeParam) {
+ int LocalInt;
+ short *LocalPointer;
+ float LocalArray[5];
+ Complete LocalComposite;
+ // Check Appertainment:
+#pragma acc host_data use_device(LocalInt)
+
+ // Valid cases:
+#pragma acc host_data use_device(LocalInt, LocalPointer, LocalArray)
+ ;
+ ;
+ // expected-error at +2{{OpenACC variable in 'use_device' clause is not a valid variable name or array name}}
+ // expected-error at +1{{OpenACC variable in 'use_device' clause is not a valid variable name or array name}}
+#pragma acc host_data use_device(LocalComposite.ScalarMember, LocalComposite.ScalarMember)
+ ;
+
+ // expected-error at +1{{OpenACC variable in 'use_device' clause is not a valid variable name or array name}}
+#pragma acc host_data use_device(LocalArray[2:1])
+
+ // expected-error at +1{{OpenACC variable in 'use_device' clause is not a valid variable name or array name}}
+#pragma acc host_data use_device(1 + IntParam)
+ ;
+
+ // expected-error at +1{{OpenACC variable in 'use_device' clause is not a valid variable name or array name}}
+#pragma acc host_data use_device(+IntParam)
+ ;
+
+ // expected-error at +2{{OpenACC sub-array length is unspecified and cannot be inferred because the subscripted value is not an array}}
+ // expected-error at +1{{OpenACC variable in 'use_device' clause is not a valid variable name or array name}}
+#pragma acc host_data use_device(PointerParam[2:])
+ ;
+
+ // expected-error at +2{{OpenACC sub-array specified range [2:5] would be out of the range of the subscripted array size of 5}}
+ // expected-error at +1{{OpenACC variable in 'use_device' clause is not a valid variable name or array name}}
+#pragma acc host_data use_device(ArrayParam[2:5])
+ ;
+
+ // expected-error at +2{{OpenACC sub-array specified range [2:5] would be out of the range of the subscripted array size of 5}}
+ // expected-error at +1{{OpenACC variable in 'use_device' clause is not a valid variable name or array name}}
+#pragma acc host_data use_device((float*)ArrayParam[2:5])
+ ;
+ // expected-error at +1{{OpenACC variable in 'use_device' clause is not a valid variable name or array name}}
+#pragma acc host_data use_device((float)ArrayParam[2])
+ ;
+
+ // expected-error at +1{{OpenACC 'use_device' clause is not valid on 'data' directive}}
+#pragma acc data use_device(LocalInt)
+ ;
+ // expected-error at +1{{OpenACC 'use_device' clause is not valid on 'enter data' directive}}
+#pragma acc enter data use_device(LocalInt)
+ // expected-error at +1{{OpenACC 'use_device' clause is not valid on 'exit data' directive}}
+#pragma acc exit data use_device(LocalInt)
+}
diff --git a/clang/test/SemaOpenACC/data-construct-wait-clause.c b/clang/test/SemaOpenACC/data-construct-wait-clause.c
index dffcba34e7b333..cef2dbdca29ed8 100644
--- a/clang/test/SemaOpenACC/data-construct-wait-clause.c
+++ b/clang/test/SemaOpenACC/data-construct-wait-clause.c
@@ -14,9 +14,8 @@ void uses() {
#pragma acc exit data copyout(arr[0]) wait(getS(), getI())
- // expected-warning at +2{{OpenACC clause 'use_device' not yet implemented}}
// expected-error at +1{{OpenACC 'wait' clause is not valid on 'host_data' directive}}
-#pragma acc host_data use_device(arr[0]) wait(getS(), getI())
+#pragma acc host_data use_device(arr) wait(getS(), getI())
;
#pragma acc data copyin(arr[0]) wait(devnum:getS(): getI())
diff --git a/clang/test/SemaOpenACC/data-construct.cpp b/clang/test/SemaOpenACC/data-construct.cpp
index 507cc5ac5cfaf5..309d0300d5a9e1 100644
--- a/clang/test/SemaOpenACC/data-construct.cpp
+++ b/clang/test/SemaOpenACC/data-construct.cpp
@@ -86,7 +86,6 @@ void AtLeastOneOf() {
#pragma acc exit data
// Host Data
- // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
#pragma acc host_data use_device(Var)
;
// OpenACC TODO: The following 'host_data' directives should diagnose, since
@@ -102,9 +101,6 @@ void AtLeastOneOf() {
void DataRules() {
int Var;
- // OpenACC TODO: Only 'async' and 'wait' are permitted after a device_type, so
- // the rest of these should diagnose.
-
// expected-error at +2{{OpenACC clause 'copy' may not follow a 'device_type' clause in a 'data' construct}}
// expected-note at +1{{previous clause is here}}
#pragma acc data device_type(*) copy(Var)
@@ -155,16 +151,13 @@ struct HasMembers {
int Member;
void HostDataError() {
- // TODO OpenACC: The following 3 should error, as use_device's var only allows
- // a variable or array, not an array index, or sub expression.
-
- // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
+ // expected-error at +1{{OpenACC variable in 'use_device' clause is not a valid variable name or array name}}
#pragma acc host_data use_device(this)
;
- // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
+ // expected-error at +1{{OpenACC variable in 'use_device' clause is not a valid variable name or array name}}
#pragma acc host_data use_device(this->Member)
;
- // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
+ // expected-error at +1{{OpenACC variable in 'use_device' clause is not a valid variable name or array name}}
#pragma acc host_data use_device(Member)
;
}
@@ -177,27 +170,22 @@ void HostDataRules() {
#pragma acc host_data if(Var) if (Var2)
;
- // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
#pragma acc host_data use_device(Var)
;
int Array[5];
- // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
#pragma acc host_data use_device(Array)
;
- // TODO OpenACC: The following 3 should error, as use_device's var only allows
- // a variable or array, not an array index, or sub expression.
-
- // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
+ // expected-error at +1{{OpenACC variable in 'use_device' clause is not a valid variable name or array name}}
#pragma acc host_data use_device(Array[1:1])
;
- // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
+ // expected-error at +1{{OpenACC variable in 'use_device' clause is not a valid variable name or array name}}
#pragma acc host_data use_device(Array[1])
;
HasMembers HM;
- // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
+ // expected-error at +1{{OpenACC variable in 'use_device' clause is not a valid variable name or array name}}
#pragma acc host_data use_device(HM.Member)
;
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 bbd04e7afa6f25..d196633c8b6d92 100644
--- a/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
+++ b/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
@@ -68,7 +68,7 @@ void uses() {
// expected-error at +1{{OpenACC 'present_or_copy' clause is not valid on 'loop' directive}}
#pragma acc loop auto present_or_copy(Var)
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
+ // expected-error at +1{{OpenACC 'use_device' clause is not valid on 'loop' directive}}
#pragma acc loop auto use_device(Var)
for(unsigned i = 0; i < 5; ++i);
// expected-error at +1{{OpenACC 'attach' clause is not valid on 'loop' directive}}
@@ -202,7 +202,7 @@ void uses() {
// expected-error at +1{{OpenACC 'present_or_copy' clause is not valid on 'loop' directive}}
#pragma acc loop present_or_copy(Var) auto
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
+ // expected-error at +1{{OpenACC 'use_device' clause is not valid on 'loop' directive}}
#pragma acc loop use_device(Var) auto
for(unsigned i = 0; i < 5; ++i);
// expected-error at +1{{OpenACC 'attach' clause is not valid on 'loop' directive}}
@@ -337,7 +337,7 @@ void uses() {
// expected-error at +1{{OpenACC 'present_or_copy' clause is not valid on 'loop' directive}}
#pragma acc loop independent present_or_copy(Var)
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
+ // expected-error at +1{{OpenACC 'use_device' clause is not valid on 'loop' directive}}
#pragma acc loop independent use_device(Var)
for(unsigned i = 0; i < 5; ++i);
// expected-error at +1{{OpenACC 'attach' clause is not valid on 'loop' directive}}
@@ -471,7 +471,7 @@ void uses() {
// expected-error at +1{{OpenACC 'present_or_copy' clause is not valid on 'loop' directive}}
#pragma acc loop present_or_copy(Var) independent
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
+ // expected-error at +1{{OpenACC 'use_device' clause is not valid on 'loop' directive}}
#pragma acc loop use_device(Var) independent
for(unsigned i = 0; i < 5; ++i);
// expected-error at +1{{OpenACC 'attach' clause is not valid on 'loop' directive}}
@@ -614,7 +614,7 @@ void uses() {
// expected-error at +1{{OpenACC 'present_or_copy' clause is not valid on 'loop' directive}}
#pragma acc loop seq present_or_copy(Var)
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
+ // expected-error at +1{{OpenACC 'use_device' clause is not valid on 'loop' directive}}
#pragma acc loop seq use_device(Var)
for(unsigned i = 0; i < 5; ++i);
// expected-error at +1{{OpenACC 'attach' clause is not valid on 'loop' directive}}
@@ -754,7 +754,7 @@ void uses() {
// expected-error at +1{{OpenACC 'present_or_copy' clause is not valid on 'loop' directive}}
#pragma acc loop present_or_copy(Var) seq
for(unsigned i = 0; i < 5; ++i);
- // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
+ // expected-error at +1{{OpenACC 'use_device' clause is not valid on 'loop' directive}}
#pragma acc loop use_device(Var) seq
for(unsigned i = 0; i < 5; ++i);
// expected-error at +1{{OpenACC 'attach' 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 8fc6690273c7de..3e4d0da60b6b27 100644
--- a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
@@ -80,8 +80,7 @@ void uses() {
// expected-error at +1{{OpenACC 'present_or_copy' clause is not valid on 'loop' directive}}
#pragma acc loop device_type(*) present_or_copy(Var)
for(int i = 0; i < 5; ++i);
- // expected-error at +2{{OpenACC clause 'use_device' may not follow a 'device_type' clause in a 'loop' construct}}
- // expected-note at +1{{previous clause is here}}
+ // expected-error at +1{{OpenACC 'use_device' clause is not valid on 'loop' directive}}
#pragma acc loop device_type(*) use_device(Var)
for(int i = 0; i < 5; ++i);
// expected-error at +1{{OpenACC 'attach' clause is not valid on 'loop' directive}}
diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
index d1a28624618990..701582138e053d 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -2894,6 +2894,11 @@ void OpenACCClauseEnqueue::VisitDeleteClause(const OpenACCDeleteClause &C) {
VisitVarList(C);
}
+void OpenACCClauseEnqueue::VisitUseDeviceClause(
+ const OpenACCUseDeviceClause &C) {
+ VisitVarList(C);
+}
+
void OpenACCClauseEnqueue::VisitDevicePtrClause(
const OpenACCDevicePtrClause &C) {
VisitVarList(C);
More information about the cfe-commits
mailing list