[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