[clang] 5d7d66b - [OpenACC] Implement 'declare' construct AST/Sema

via cfe-commits cfe-commits at lists.llvm.org
Mon Mar 3 07:48:35 PST 2025


Author: erichkeane
Date: 2025-03-03T07:48:29-08:00
New Revision: 5d7d66ba0d1ad6fcf1aefffd045eea88597f4614

URL: https://github.com/llvm/llvm-project/commit/5d7d66ba0d1ad6fcf1aefffd045eea88597f4614
DIFF: https://github.com/llvm/llvm-project/commit/5d7d66ba0d1ad6fcf1aefffd045eea88597f4614.diff

LOG: [OpenACC] Implement 'declare' construct AST/Sema

The 'declare' construct is the first of two 'declaration' level
constructs, so it is legal in any place a declaration is, including as a
statement, which this accomplishes by wrapping it in a DeclStmt. All
clauses on this have a 'same scope' requirement, which this enforces as
declaration context instead, which makes it possible to implement these
as a template.

The 'link' and 'device_resident' clauses are also added, which have some
similar/small restrictions, but are otherwise pretty rote.

This patch implements all of the above.

Added: 
    clang/include/clang/AST/DeclOpenACC.h
    clang/lib/AST/DeclOpenACC.cpp
    clang/test/AST/ast-print-openacc-declare-construct.cpp
    clang/test/SemaOpenACC/declare-construct-ast.cpp
    clang/test/SemaOpenACC/declare-construct.cpp

Modified: 
    clang/include/clang/AST/DeclVisitor.h
    clang/include/clang/AST/JSONNodeDumper.h
    clang/include/clang/AST/OpenACCClause.h
    clang/include/clang/AST/RecursiveASTVisitor.h
    clang/include/clang/AST/TextNodeDumper.h
    clang/include/clang/Basic/DeclNodes.td
    clang/include/clang/Basic/DiagnosticSemaKinds.td
    clang/include/clang/Basic/OpenACCClauses.def
    clang/include/clang/Parse/Parser.h
    clang/include/clang/Sema/SemaBase.h
    clang/include/clang/Sema/SemaOpenACC.h
    clang/include/clang/Serialization/ASTBitCodes.h
    clang/include/clang/Serialization/ASTRecordReader.h
    clang/lib/AST/ASTStructuralEquivalence.cpp
    clang/lib/AST/CMakeLists.txt
    clang/lib/AST/DeclBase.cpp
    clang/lib/AST/DeclPrinter.cpp
    clang/lib/AST/JSONNodeDumper.cpp
    clang/lib/AST/OpenACCClause.cpp
    clang/lib/AST/StmtPrinter.cpp
    clang/lib/AST/StmtProfile.cpp
    clang/lib/AST/TextNodeDumper.cpp
    clang/lib/CodeGen/CGDecl.cpp
    clang/lib/CodeGen/CodeGenModule.h
    clang/lib/Parse/ParseOpenACC.cpp
    clang/lib/Sema/SemaBase.cpp
    clang/lib/Sema/SemaOpenACC.cpp
    clang/lib/Sema/SemaOpenACCClause.cpp
    clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
    clang/lib/Sema/TreeTransform.h
    clang/lib/Serialization/ASTCommon.cpp
    clang/lib/Serialization/ASTReader.cpp
    clang/lib/Serialization/ASTReaderDecl.cpp
    clang/lib/Serialization/ASTWriter.cpp
    clang/lib/Serialization/ASTWriterDecl.cpp
    clang/test/ParserOpenACC/parse-clauses.c
    clang/test/ParserOpenACC/parse-constructs.c
    clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
    clang/test/SemaOpenACC/combined-construct-device_type-clause.c
    clang/test/SemaOpenACC/compute-construct-device_type-clause.c
    clang/test/SemaOpenACC/data-construct.cpp
    clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
    clang/test/SemaOpenACC/loop-construct-device_type-clause.c
    clang/test/SemaOpenACC/unimplemented-construct.c
    clang/tools/libclang/CIndex.cpp

Removed: 
    


################################################################################
diff  --git a/clang/include/clang/AST/DeclOpenACC.h b/clang/include/clang/AST/DeclOpenACC.h
new file mode 100644
index 0000000000000..fdeb1315cfd09
--- /dev/null
+++ b/clang/include/clang/AST/DeclOpenACC.h
@@ -0,0 +1,106 @@
+//=- DeclOpenACC.h - Classes for representing OpenACC directives -*- C++ -*-==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// This file defines OpenACC nodes for declarative directives.
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_CLANG_AST_DECLOPENACC_H
+#define LLVM_CLANG_AST_DECLOPENACC_H
+
+#include "clang/AST/ASTContext.h"
+#include "clang/AST/Decl.h"
+#include "clang/AST/OpenACCClause.h"
+#include "clang/Basic/OpenACCKinds.h"
+
+namespace clang {
+
+// A base class for the declaration constructs, which manages the clauses and
+// basic source location information. Currently not part of the Decl inheritence
+// tree, as we should never have a reason to store one of these.
+class OpenACCConstructDecl : public Decl {
+  friend class ASTDeclReader;
+  friend class ASTDeclWriter;
+  // The directive kind, each implementation of this interface is expected to
+  // handle a specific kind.
+  OpenACCDirectiveKind DirKind = OpenACCDirectiveKind::Invalid;
+  SourceLocation DirectiveLoc;
+  SourceLocation EndLoc;
+  /// The list of clauses.  This is stored here as an ArrayRef, as this is the
+  /// most convienient place to access the list, however the list itself should
+  /// be stored in leaf nodes, likely in trailing-storage.
+  MutableArrayRef<const OpenACCClause *> Clauses;
+
+protected:
+  OpenACCConstructDecl(Kind DeclKind, DeclContext *DC, OpenACCDirectiveKind K,
+                       SourceLocation StartLoc, SourceLocation DirLoc,
+                       SourceLocation EndLoc)
+      : Decl(DeclKind, DC, StartLoc), DirKind(K), DirectiveLoc(DirLoc),
+        EndLoc(EndLoc) {}
+
+  OpenACCConstructDecl(Kind DeclKind) : Decl(DeclKind, EmptyShell{}) {}
+
+  void setClauseList(MutableArrayRef<const OpenACCClause *> NewClauses) {
+    assert(Clauses.empty() && "Cannot change clause list");
+    Clauses = NewClauses;
+  }
+
+public:
+  OpenACCDirectiveKind getDirectiveKind() const { return DirKind; }
+  SourceLocation getDirectiveLoc() const { return DirectiveLoc; }
+  virtual SourceRange getSourceRange() const override LLVM_READONLY {
+    return SourceRange(getLocation(), EndLoc);
+  }
+
+  ArrayRef<const OpenACCClause *> clauses() const { return Clauses; }
+};
+
+class OpenACCDeclareDecl final
+    : public OpenACCConstructDecl,
+      private llvm::TrailingObjects<OpenACCDeclareDecl, const OpenACCClause *> {
+  friend TrailingObjects;
+  friend class ASTDeclReader;
+  friend class ASTDeclWriter;
+
+  OpenACCDeclareDecl(unsigned NumClauses)
+      : OpenACCConstructDecl(OpenACCDeclare) {
+    std::uninitialized_value_construct(
+        getTrailingObjects<const OpenACCClause *>(),
+        getTrailingObjects<const OpenACCClause *>() + NumClauses);
+    setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
+                                  NumClauses));
+  }
+
+  OpenACCDeclareDecl(DeclContext *DC, SourceLocation StartLoc,
+                     SourceLocation DirLoc, SourceLocation EndLoc,
+                     ArrayRef<const OpenACCClause *> Clauses)
+      : OpenACCConstructDecl(OpenACCDeclare, DC, OpenACCDirectiveKind::Declare,
+                             StartLoc, DirLoc, EndLoc) {
+    // Initialize the trailing storage.
+    std::uninitialized_copy(Clauses.begin(), Clauses.end(),
+                            getTrailingObjects<const OpenACCClause *>());
+
+    setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
+                                  Clauses.size()));
+  }
+
+public:
+  static OpenACCDeclareDecl *Create(ASTContext &Ctx, DeclContext *DC,
+                                    SourceLocation StartLoc,
+                                    SourceLocation DirLoc,
+                                    SourceLocation EndLoc,
+                                    ArrayRef<const OpenACCClause *> Clauses);
+  static OpenACCDeclareDecl *
+  CreateDeserialized(ASTContext &Ctx, GlobalDeclID ID, unsigned NumClauses);
+  static bool classof(const Decl *D) { return classofKind(D->getKind()); }
+  static bool classofKind(Kind K) { return K == OpenACCDeclare; }
+};
+} // namespace clang
+
+#endif

diff  --git a/clang/include/clang/AST/DeclVisitor.h b/clang/include/clang/AST/DeclVisitor.h
index 8690cdda4bb70..8b7c30ee4051e 100644
--- a/clang/include/clang/AST/DeclVisitor.h
+++ b/clang/include/clang/AST/DeclVisitor.h
@@ -18,6 +18,7 @@
 #include "clang/AST/DeclCXX.h"
 #include "clang/AST/DeclFriend.h"
 #include "clang/AST/DeclObjC.h"
+#include "clang/AST/DeclOpenACC.h"
 #include "clang/AST/DeclOpenMP.h"
 #include "clang/AST/DeclTemplate.h"
 #include "llvm/ADT/STLExtras.h"

diff  --git a/clang/include/clang/AST/JSONNodeDumper.h b/clang/include/clang/AST/JSONNodeDumper.h
index 660a05435003d..97b07e1afbf68 100644
--- a/clang/include/clang/AST/JSONNodeDumper.h
+++ b/clang/include/clang/AST/JSONNodeDumper.h
@@ -281,6 +281,8 @@ class JSONNodeDumper
   void VisitObjCPropertyImplDecl(const ObjCPropertyImplDecl *D);
   void VisitBlockDecl(const BlockDecl *D);
 
+  void VisitOpenACCDeclareDecl(const OpenACCDeclareDecl *D);
+
   void VisitDeclRefExpr(const DeclRefExpr *DRE);
   void VisitSYCLUniqueStableNameExpr(const SYCLUniqueStableNameExpr *E);
   void VisitOpenACCAsteriskSizeExpr(const OpenACCAsteriskSizeExpr *E);

diff  --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h
index f5be54bdada8b..0fc638b73bbd2 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -1171,6 +1171,55 @@ class OpenACCReductionClause final
   OpenACCReductionOperator getReductionOp() const { return Op; }
 };
 
+class OpenACCLinkClause final
+    : public OpenACCClauseWithVarList,
+      private llvm::TrailingObjects<OpenACCLinkClause, Expr *> {
+  friend TrailingObjects;
+
+  OpenACCLinkClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
+                    ArrayRef<Expr *> VarList, SourceLocation EndLoc)
+      : OpenACCClauseWithVarList(OpenACCClauseKind::Link, 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::Link;
+  }
+
+  static OpenACCLinkClause *Create(const ASTContext &C, SourceLocation BeginLoc,
+                                   SourceLocation LParenLoc,
+                                   ArrayRef<Expr *> VarList,
+                                   SourceLocation EndLoc);
+};
+
+class OpenACCDeviceResidentClause final
+    : public OpenACCClauseWithVarList,
+      private llvm::TrailingObjects<OpenACCDeviceResidentClause, Expr *> {
+  friend TrailingObjects;
+
+  OpenACCDeviceResidentClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
+                              ArrayRef<Expr *> VarList, SourceLocation EndLoc)
+      : OpenACCClauseWithVarList(OpenACCClauseKind::DeviceResident, 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::DeviceResident;
+  }
+
+  static OpenACCDeviceResidentClause *
+  Create(const ASTContext &C, SourceLocation BeginLoc, SourceLocation LParenLoc,
+         ArrayRef<Expr *> VarList, SourceLocation EndLoc);
+};
+
 template <class Impl> class OpenACCClauseVisitor {
   Impl &getDerived() { return static_cast<Impl &>(*this); }
 

diff  --git a/clang/include/clang/AST/RecursiveASTVisitor.h b/clang/include/clang/AST/RecursiveASTVisitor.h
index 5964cbaec8e44..b1b4363b65721 100644
--- a/clang/include/clang/AST/RecursiveASTVisitor.h
+++ b/clang/include/clang/AST/RecursiveASTVisitor.h
@@ -20,6 +20,7 @@
 #include "clang/AST/DeclCXX.h"
 #include "clang/AST/DeclFriend.h"
 #include "clang/AST/DeclObjC.h"
+#include "clang/AST/DeclOpenACC.h"
 #include "clang/AST/DeclOpenMP.h"
 #include "clang/AST/DeclTemplate.h"
 #include "clang/AST/DeclarationName.h"
@@ -1821,6 +1822,9 @@ DEF_TRAVERSE_DECL(OMPAllocateDecl, {
     TRY_TO(TraverseOMPClause(C));
 })
 
+DEF_TRAVERSE_DECL(OpenACCDeclareDecl,
+                  { TRY_TO(VisitOpenACCClauseList(D->clauses())); })
+
 // A helper method for TemplateDecl's children.
 template <typename Derived>
 bool RecursiveASTVisitor<Derived>::TraverseTemplateParameterListHelper(

diff  --git a/clang/include/clang/AST/TextNodeDumper.h b/clang/include/clang/AST/TextNodeDumper.h
index 81844db2c77fa..b0e5b28e22863 100644
--- a/clang/include/clang/AST/TextNodeDumper.h
+++ b/clang/include/clang/AST/TextNodeDumper.h
@@ -423,6 +423,7 @@ class TextNodeDumper
   void VisitOpenACCUpdateConstruct(const OpenACCUpdateConstruct *S);
   void VisitOpenACCAtomicConstruct(const OpenACCAtomicConstruct *S);
   void VisitOpenACCAsteriskSizeExpr(const OpenACCAsteriskSizeExpr *S);
+  void VisitOpenACCDeclareDecl(const OpenACCDeclareDecl *D);
   void VisitEmbedExpr(const EmbedExpr *S);
   void VisitAtomicExpr(const AtomicExpr *AE);
   void VisitConvertVectorExpr(const ConvertVectorExpr *S);

diff  --git a/clang/include/clang/Basic/DeclNodes.td b/clang/include/clang/Basic/DeclNodes.td
index 723113dc2486e..227184dc71cc7 100644
--- a/clang/include/clang/Basic/DeclNodes.td
+++ b/clang/include/clang/Basic/DeclNodes.td
@@ -111,3 +111,4 @@ def Empty : DeclNode<Decl>;
 def RequiresExprBody : DeclNode<Decl>, DeclContext;
 def LifetimeExtendedTemporary : DeclNode<Decl>;
 def HLSLBuffer : DeclNode<Named, "HLSLBuffer">, DeclContext;
+def OpenACCDeclare : DeclNode<Decl, "#pragma acc declare">;

diff  --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index d094c075ecee2..0efb15405ed5d 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12812,6 +12812,7 @@ def err_acc_duplicate_clause_disallowed
             "directive">;
 def note_acc_previous_clause_here : Note<"previous clause is here">;
 def note_acc_previous_expr_here : Note<"previous expression is here">;
+def note_acc_previous_reference : Note<"previous reference is here">;
 def err_acc_branch_in_out_compute_construct
     : Error<"invalid %select{branch|return|throw}0 %select{out of|into}1 "
             "OpenACC Compute/Combined Construct">;
@@ -12844,9 +12845,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_not_a_var_ref_use_device_declare
+    : Error<"OpenACC variable %select{in 'use_device' clause|on 'declare' "
+            "construct}0 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
@@ -13014,6 +13015,23 @@ def note_acc_atomic_mismatch_compound_operand
            "side of assignment|<not possible>|on left hand side of compound "
            "assignment|on left hand side of assignment}0('%1') from the first "
            "statement">;
+def err_acc_declare_required_clauses
+    : Error<"no valid clauses specified in OpenACC 'declare' directive">;
+def err_acc_declare_clause_at_global
+    : Error<"OpenACC '%0' clause on a 'declare' directive is not allowed at "
+            "global or namespace scope">;
+def err_acc_link_not_extern
+    : Error<"variable referenced by 'link' clause not in global or namespace "
+            "scope must be marked 'extern'">;
+def err_acc_declare_extern
+    : Error<"'extern' variable may not be referenced by '%0' clause on an "
+            "OpenACC 'declare' directive">;
+def err_acc_declare_same_scope
+    : Error<"variable appearing in '%0' clause of OpenACC 'declare' directive "
+            "must be in the same scope as the directive">;
+def err_acc_multiple_references
+    : Error<"variable referenced in '%0' clause of OpenACC 'declare' directive "
+            "was already referenced">;
 
 // AMDGCN builtins diagnostics
 def err_amdgcn_global_load_lds_size_invalid_value : Error<"invalid size value">;

diff  --git a/clang/include/clang/Basic/OpenACCClauses.def b/clang/include/clang/Basic/OpenACCClauses.def
index 8b15007c85557..1d9e3b09a749b 100644
--- a/clang/include/clang/Basic/OpenACCClauses.def
+++ b/clang/include/clang/Basic/OpenACCClauses.def
@@ -44,6 +44,7 @@ VISIT_CLAUSE(Detach)
 VISIT_CLAUSE(Device)
 VISIT_CLAUSE(DeviceNum)
 VISIT_CLAUSE(DevicePtr)
+VISIT_CLAUSE(DeviceResident)
 VISIT_CLAUSE(DeviceType)
 CLAUSE_ALIAS(DType, DeviceType, false)
 VISIT_CLAUSE(Finalize)
@@ -53,6 +54,7 @@ VISIT_CLAUSE(Host)
 VISIT_CLAUSE(If)
 VISIT_CLAUSE(IfPresent)
 VISIT_CLAUSE(Independent)
+VISIT_CLAUSE(Link)
 VISIT_CLAUSE(NoCreate)
 VISIT_CLAUSE(NumGangs)
 VISIT_CLAUSE(NumWorkers)

diff  --git a/clang/include/clang/Parse/Parser.h b/clang/include/clang/Parse/Parser.h
index 335258d597028..0602f44333b20 100644
--- a/clang/include/clang/Parse/Parser.h
+++ b/clang/include/clang/Parse/Parser.h
@@ -3756,9 +3756,11 @@ class Parser : public CodeCompletionHandler {
 
   using OpenACCVarParseResult = std::pair<ExprResult, OpenACCParseCanContinue>;
   /// Parses a single variable in a variable list for OpenACC.
-  OpenACCVarParseResult ParseOpenACCVar(OpenACCClauseKind CK);
+  OpenACCVarParseResult ParseOpenACCVar(OpenACCDirectiveKind DK,
+                                        OpenACCClauseKind CK);
   /// Parses the variable list for the variety of places that take a var-list.
-  llvm::SmallVector<Expr *> ParseOpenACCVarList(OpenACCClauseKind CK);
+  llvm::SmallVector<Expr *> ParseOpenACCVarList(OpenACCDirectiveKind DK,
+                                                OpenACCClauseKind CK);
   /// Parses any parameters for an OpenACC Clause, including required/optional
   /// parens.
   OpenACCClauseParseResult

diff  --git a/clang/include/clang/Sema/SemaBase.h b/clang/include/clang/Sema/SemaBase.h
index 0b05245ab9686..463cae83c7e81 100644
--- a/clang/include/clang/Sema/SemaBase.h
+++ b/clang/include/clang/Sema/SemaBase.h
@@ -42,6 +42,7 @@ class SemaBase {
   ASTContext &getASTContext() const;
   DiagnosticsEngine &getDiagnostics() const;
   const LangOptions &getLangOpts() const;
+  DeclContext *getCurContext() const;
 
   /// Helper class that creates diagnostics with optional
   /// template instantiation stacks.

diff  --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h
index 3004b98760a98..36a2b6ff7edc3 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -33,6 +33,9 @@ class IdentifierInfo;
 class OpenACCClause;
 
 class SemaOpenACC : public SemaBase {
+public:
+  using DeclGroupPtrTy = OpaquePtr<DeclGroupRef>;
+
 private:
   struct ComputeConstructInfo {
     /// Which type of compute construct we are inside of, which we can use to
@@ -158,6 +161,13 @@ class SemaOpenACC : public SemaBase {
   /// Helper function for checking the 'for' and 'range for' stmts.
   void ForStmtBeginHelper(SourceLocation ForLoc, ForStmtBeginChecker &C);
 
+  // The 'declare' construct requires only a single reference among ALL declare
+  // directives in a context. We store existing references to check. Because the
+  // rules prevent referencing the same variable from multiple declaration
+  // contexts, we can just store the declaration and location of the reference.
+  llvm::DenseMap<const clang::DeclaratorDecl *, SourceLocation>
+      DeclareVarReferences;
+
 public:
   ComputeConstructInfo &getActiveComputeConstructInfo() {
     return ActiveComputeConstructInfo;
@@ -411,6 +421,8 @@ class SemaOpenACC : public SemaBase {
               ClauseKind == OpenACCClauseKind::Reduction ||
               ClauseKind == OpenACCClauseKind::Host ||
               ClauseKind == OpenACCClauseKind::Device ||
+              ClauseKind == OpenACCClauseKind::DeviceResident ||
+              ClauseKind == OpenACCClauseKind::Link ||
               (ClauseKind == OpenACCClauseKind::Self &&
                DirKind == OpenACCDirectiveKind::Update) ||
               ClauseKind == OpenACCClauseKind::FirstPrivate) &&
@@ -427,23 +439,10 @@ class SemaOpenACC : public SemaBase {
     }
 
     bool isReadOnly() const {
-      assert((ClauseKind == OpenACCClauseKind::CopyIn ||
-              ClauseKind == OpenACCClauseKind::PCopyIn ||
-              ClauseKind == OpenACCClauseKind::PresentOrCopyIn) &&
-             "Only copyin accepts 'readonly:' tag");
       return std::get<VarListDetails>(Details).IsReadOnly;
     }
 
-    bool isZero() const {
-      assert((ClauseKind == OpenACCClauseKind::CopyOut ||
-              ClauseKind == OpenACCClauseKind::PCopyOut ||
-              ClauseKind == OpenACCClauseKind::PresentOrCopyOut ||
-              ClauseKind == OpenACCClauseKind::Create ||
-              ClauseKind == OpenACCClauseKind::PCreate ||
-              ClauseKind == OpenACCClauseKind::PresentOrCreate) &&
-             "Only copyout/create accepts 'zero' tag");
-      return std::get<VarListDetails>(Details).IsZero;
-    }
+    bool isZero() const { return std::get<VarListDetails>(Details).IsZero; }
 
     bool isForce() const {
       assert(ClauseKind == OpenACCClauseKind::Collapse &&
@@ -557,6 +556,8 @@ class SemaOpenACC : public SemaBase {
               ClauseKind == OpenACCClauseKind::DevicePtr ||
               ClauseKind == OpenACCClauseKind::Host ||
               ClauseKind == OpenACCClauseKind::Device ||
+              ClauseKind == OpenACCClauseKind::DeviceResident ||
+              ClauseKind == OpenACCClauseKind::Link ||
               (ClauseKind == OpenACCClauseKind::Self &&
                DirKind == OpenACCDirectiveKind::Update) ||
               ClauseKind == OpenACCClauseKind::FirstPrivate) &&
@@ -600,6 +601,8 @@ class SemaOpenACC : public SemaBase {
               ClauseKind == OpenACCClauseKind::DevicePtr ||
               ClauseKind == OpenACCClauseKind::Host ||
               ClauseKind == OpenACCClauseKind::Device ||
+              ClauseKind == OpenACCClauseKind::DeviceResident ||
+              ClauseKind == OpenACCClauseKind::Link ||
               (ClauseKind == OpenACCClauseKind::Self &&
                DirKind == OpenACCDirectiveKind::Update) ||
               ClauseKind == OpenACCClauseKind::FirstPrivate) &&
@@ -744,7 +747,11 @@ class SemaOpenACC : public SemaBase {
 
   /// Called after the directive has been completely parsed, including the
   /// declaration group or associated statement.
-  DeclGroupRef ActOnEndDeclDirective();
+  DeclGroupRef ActOnEndDeclDirective(OpenACCDirectiveKind K,
+                                     SourceLocation StartLoc,
+                                     SourceLocation DirLoc,
+                                     SourceLocation EndLoc,
+                                     ArrayRef<OpenACCClause *> Clauses);
 
   /// Called when encountering an 'int-expr' for OpenACC, and manages
   /// conversions and diagnostics to 'int'.
@@ -753,7 +760,16 @@ class SemaOpenACC : public SemaBase {
 
   /// Called when encountering a 'var' for OpenACC, ensures it is actually a
   /// declaration reference to a variable of the correct type.
-  ExprResult ActOnVar(OpenACCClauseKind CK, Expr *VarExpr);
+  ExprResult ActOnVar(OpenACCDirectiveKind DK, OpenACCClauseKind CK,
+                      Expr *VarExpr);
+
+  // Called after 'ActOnVar' specifically for a 'link' clause, which has to do
+  // some minor additional checks.
+  llvm::SmallVector<Expr *> CheckLinkClauseVarList(ArrayRef<Expr *> VarExpr);
+
+  // Checking for the arguments specific to the declare-clause that need to be
+  // checked during both phases of template translation.
+  bool CheckDeclareClause(SemaOpenACC::OpenACCParsedClause &Clause);
 
   /// Called while semantically analyzing the reduction clause, ensuring the var
   /// is the correct kind of reference.

diff  --git a/clang/include/clang/Serialization/ASTBitCodes.h b/clang/include/clang/Serialization/ASTBitCodes.h
index c998be34b9d89..5698bebf13445 100644
--- a/clang/include/clang/Serialization/ASTBitCodes.h
+++ b/clang/include/clang/Serialization/ASTBitCodes.h
@@ -1528,7 +1528,10 @@ enum DeclCode {
   // A decls specialization record.
   DECL_PARTIAL_SPECIALIZATIONS,
 
-  DECL_LAST = DECL_PARTIAL_SPECIALIZATIONS
+  // An OpenACCDeclareDecl record.
+  DECL_OPENACC_DECLARE,
+
+  DECL_LAST = DECL_OPENACC_DECLARE
 };
 
 /// Record codes for each kind of statement or expression.

diff  --git a/clang/include/clang/Serialization/ASTRecordReader.h b/clang/include/clang/Serialization/ASTRecordReader.h
index 2561418b78ca7..ae7cd84fbc647 100644
--- a/clang/include/clang/Serialization/ASTRecordReader.h
+++ b/clang/include/clang/Serialization/ASTRecordReader.h
@@ -278,7 +278,8 @@ class ASTRecordReader
   /// Read an OpenACC clause, advancing Idx.
   OpenACCClause *readOpenACCClause();
 
-  /// Read a list of OpenACC clauses into the passed SmallVector.
+  /// Read a list of OpenACC clauses into the passed SmallVector, during
+  /// statement reading.
   void readOpenACCClauseList(MutableArrayRef<const OpenACCClause *> Clauses);
 
   /// Read a source location, advancing Idx.

diff  --git a/clang/lib/AST/ASTStructuralEquivalence.cpp b/clang/lib/AST/ASTStructuralEquivalence.cpp
index eaf0748395268..b74f67f0a9fed 100644
--- a/clang/lib/AST/ASTStructuralEquivalence.cpp
+++ b/clang/lib/AST/ASTStructuralEquivalence.cpp
@@ -66,6 +66,7 @@
 #include "clang/AST/DeclCXX.h"
 #include "clang/AST/DeclFriend.h"
 #include "clang/AST/DeclObjC.h"
+#include "clang/AST/DeclOpenACC.h"
 #include "clang/AST/DeclOpenMP.h"
 #include "clang/AST/DeclTemplate.h"
 #include "clang/AST/ExprCXX.h"

diff  --git a/clang/lib/AST/CMakeLists.txt b/clang/lib/AST/CMakeLists.txt
index cb13c5225b713..f6056e3935a63 100644
--- a/clang/lib/AST/CMakeLists.txt
+++ b/clang/lib/AST/CMakeLists.txt
@@ -50,6 +50,7 @@ add_clang_library(clangAST
   DeclFriend.cpp
   DeclGroup.cpp
   DeclObjC.cpp
+  DeclOpenACC.cpp
   DeclOpenMP.cpp
   DeclPrinter.cpp
   DeclTemplate.cpp

diff  --git a/clang/lib/AST/DeclBase.cpp b/clang/lib/AST/DeclBase.cpp
index ab9d4869a74ee..9d8eb07f255ed 100644
--- a/clang/lib/AST/DeclBase.cpp
+++ b/clang/lib/AST/DeclBase.cpp
@@ -21,6 +21,7 @@
 #include "clang/AST/DeclContextInternals.h"
 #include "clang/AST/DeclFriend.h"
 #include "clang/AST/DeclObjC.h"
+#include "clang/AST/DeclOpenACC.h"
 #include "clang/AST/DeclOpenMP.h"
 #include "clang/AST/DeclTemplate.h"
 #include "clang/AST/DependentDiagnostic.h"
@@ -992,6 +993,7 @@ unsigned Decl::getIdentifierNamespaceForKind(Kind DeclKind) {
     case LifetimeExtendedTemporary:
     case RequiresExprBody:
     case ImplicitConceptSpecialization:
+    case OpenACCDeclare:
       // Never looked up by name.
       return 0;
   }

diff  --git a/clang/lib/AST/DeclOpenACC.cpp b/clang/lib/AST/DeclOpenACC.cpp
new file mode 100644
index 0000000000000..568fd9dcceec7
--- /dev/null
+++ b/clang/lib/AST/DeclOpenACC.cpp
@@ -0,0 +1,33 @@
+//===--- DeclOpenACC.cpp - Classes for OpenACC Constructs -----------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file implements the subclasses of Decl class declared in Decl.h
+//
+//===----------------------------------------------------------------------===//
+
+#include "clang/AST/DeclOpenACC.h"
+#include "clang/AST/ASTContext.h"
+
+using namespace clang;
+
+OpenACCDeclareDecl *
+OpenACCDeclareDecl::Create(ASTContext &Ctx, DeclContext *DC,
+                           SourceLocation StartLoc, SourceLocation DirLoc,
+                           SourceLocation EndLoc,
+                           ArrayRef<const OpenACCClause *> Clauses) {
+  return new (Ctx, DC,
+              additionalSizeToAlloc<const OpenACCClause *>(Clauses.size()))
+      OpenACCDeclareDecl(DC, StartLoc, DirLoc, EndLoc, Clauses);
+}
+
+OpenACCDeclareDecl *
+OpenACCDeclareDecl::CreateDeserialized(ASTContext &Ctx, GlobalDeclID ID,
+                                       unsigned NumClauses) {
+  return new (Ctx, ID, additionalSizeToAlloc<const OpenACCClause *>(NumClauses))
+      OpenACCDeclareDecl(NumClauses);
+}

diff  --git a/clang/lib/AST/DeclPrinter.cpp b/clang/lib/AST/DeclPrinter.cpp
index 0d51fdbc7e126..5c299567ce567 100644
--- a/clang/lib/AST/DeclPrinter.cpp
+++ b/clang/lib/AST/DeclPrinter.cpp
@@ -112,6 +112,8 @@ namespace {
     void VisitNonTypeTemplateParmDecl(const NonTypeTemplateParmDecl *NTTP);
     void VisitHLSLBufferDecl(HLSLBufferDecl *D);
 
+    void VisitOpenACCDeclareDecl(OpenACCDeclareDecl *D);
+
     void printTemplateParameters(const TemplateParameterList *Params,
                                  bool OmitTemplateKW = false);
     void printTemplateArguments(llvm::ArrayRef<TemplateArgument> Args,
@@ -495,6 +497,8 @@ void DeclPrinter::VisitDeclContext(DeclContext *DC, bool Indent) {
         isa<OMPDeclareMapperDecl>(*D) || isa<OMPRequiresDecl>(*D) ||
         isa<OMPAllocateDecl>(*D))
       Terminator = nullptr;
+    else if (isa<OpenACCDeclareDecl>(*D))
+      Terminator = nullptr;
     else if (isa<ObjCMethodDecl>(*D) && cast<ObjCMethodDecl>(*D)->hasBody())
       Terminator = nullptr;
     else if (auto FD = dyn_cast<FunctionDecl>(*D)) {
@@ -1910,3 +1914,11 @@ void DeclPrinter::VisitNonTypeTemplateParmDecl(
                                                    /*IncludeType=*/false);
   }
 }
+
+void DeclPrinter::VisitOpenACCDeclareDecl(OpenACCDeclareDecl *D) {
+  if (!D->isInvalidDecl()) {
+    Out << "#pragma acc declare ";
+    OpenACCClausePrinter Printer(Out, Policy);
+    Printer.VisitClauseList(D->clauses());
+  }
+}

diff  --git a/clang/lib/AST/JSONNodeDumper.cpp b/clang/lib/AST/JSONNodeDumper.cpp
index 27fd214dcee3b..d692033830191 100644
--- a/clang/lib/AST/JSONNodeDumper.cpp
+++ b/clang/lib/AST/JSONNodeDumper.cpp
@@ -1362,6 +1362,8 @@ void JSONNodeDumper::VisitSYCLUniqueStableNameExpr(
 void JSONNodeDumper::VisitOpenACCAsteriskSizeExpr(
     const OpenACCAsteriskSizeExpr *E) {}
 
+void JSONNodeDumper::VisitOpenACCDeclareDecl(const OpenACCDeclareDecl *D) {}
+
 void JSONNodeDumper::VisitPredefinedExpr(const PredefinedExpr *PE) {
   JOS.attribute("name", PredefinedExpr::getIdentKindName(PE->getIdentKind()));
 }

diff  --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp
index aa14ab902ba66..cef241eface91 100644
--- a/clang/lib/AST/OpenACCClause.cpp
+++ b/clang/lib/AST/OpenACCClause.cpp
@@ -40,6 +40,8 @@ bool OpenACCClauseWithVarList::classof(const OpenACCClause *C) {
          OpenACCCopyInClause::classof(C) || OpenACCCopyOutClause::classof(C) ||
          OpenACCReductionClause::classof(C) ||
          OpenACCCreateClause::classof(C) || OpenACCDeviceClause::classof(C) ||
+         OpenACCLinkClause::classof(C) ||
+         OpenACCDeviceResidentClause::classof(C) ||
          OpenACCHostClause::classof(C);
 }
 bool OpenACCClauseWithCondition::classof(const OpenACCClause *C) {
@@ -438,6 +440,25 @@ OpenACCCopyClause::Create(const ASTContext &C, OpenACCClauseKind Spelling,
       OpenACCCopyClause(Spelling, BeginLoc, LParenLoc, VarList, EndLoc);
 }
 
+OpenACCLinkClause *OpenACCLinkClause::Create(const ASTContext &C,
+                                             SourceLocation BeginLoc,
+                                             SourceLocation LParenLoc,
+                                             ArrayRef<Expr *> VarList,
+                                             SourceLocation EndLoc) {
+  void *Mem =
+      C.Allocate(OpenACCLinkClause::totalSizeToAlloc<Expr *>(VarList.size()));
+  return new (Mem) OpenACCLinkClause(BeginLoc, LParenLoc, VarList, EndLoc);
+}
+
+OpenACCDeviceResidentClause *OpenACCDeviceResidentClause::Create(
+    const ASTContext &C, SourceLocation BeginLoc, SourceLocation LParenLoc,
+    ArrayRef<Expr *> VarList, SourceLocation EndLoc) {
+  void *Mem = C.Allocate(
+      OpenACCDeviceResidentClause::totalSizeToAlloc<Expr *>(VarList.size()));
+  return new (Mem)
+      OpenACCDeviceResidentClause(BeginLoc, LParenLoc, VarList, EndLoc);
+}
+
 OpenACCCopyInClause *
 OpenACCCopyInClause::Create(const ASTContext &C, OpenACCClauseKind Spelling,
                             SourceLocation BeginLoc, SourceLocation LParenLoc,
@@ -754,6 +775,21 @@ void OpenACCClausePrinter::VisitCopyClause(const OpenACCCopyClause &C) {
   OS << ")";
 }
 
+void OpenACCClausePrinter::VisitLinkClause(const OpenACCLinkClause &C) {
+  OS << "link(";
+  llvm::interleaveComma(C.getVarList(), OS,
+                        [&](const Expr *E) { printExpr(E); });
+  OS << ")";
+}
+
+void OpenACCClausePrinter::VisitDeviceResidentClause(
+    const OpenACCDeviceResidentClause &C) {
+  OS << "device_resident(";
+  llvm::interleaveComma(C.getVarList(), OS,
+                        [&](const Expr *E) { printExpr(E); });
+  OS << ")";
+}
+
 void OpenACCClausePrinter::VisitCopyInClause(const OpenACCCopyInClause &C) {
   OS << C.getClauseKind() << '(';
   if (C.isReadOnly())

diff  --git a/clang/lib/AST/StmtPrinter.cpp b/clang/lib/AST/StmtPrinter.cpp
index c8ea7b52a6241..dac9a95e10f3d 100644
--- a/clang/lib/AST/StmtPrinter.cpp
+++ b/clang/lib/AST/StmtPrinter.cpp
@@ -17,6 +17,7 @@
 #include "clang/AST/DeclBase.h"
 #include "clang/AST/DeclCXX.h"
 #include "clang/AST/DeclObjC.h"
+#include "clang/AST/DeclOpenACC.h"
 #include "clang/AST/DeclOpenMP.h"
 #include "clang/AST/DeclTemplate.h"
 #include "clang/AST/Expr.h"
@@ -262,7 +263,10 @@ void StmtPrinter::VisitNullStmt(NullStmt *Node) {
 void StmtPrinter::VisitDeclStmt(DeclStmt *Node) {
   Indent();
   PrintRawDeclStmt(Node);
-  OS << ";" << NL;
+  // Certain pragma declarations shouldn't have a semi-colon after them.
+  if (!Node->isSingleDecl() || !isa<OpenACCDeclareDecl>(Node->getSingleDecl()))
+    OS << ";";
+  OS << NL;
 }
 
 void StmtPrinter::VisitCompoundStmt(CompoundStmt *Node) {

diff  --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index 2603df25ba2a4..e283a9ad4a567 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -2550,6 +2550,16 @@ void OpenACCClauseProfiler::VisitIfClause(const OpenACCIfClause &Clause) {
 void OpenACCClauseProfiler::VisitCopyClause(const OpenACCCopyClause &Clause) {
   VisitClauseWithVarList(Clause);
 }
+
+void OpenACCClauseProfiler::VisitLinkClause(const OpenACCLinkClause &Clause) {
+  VisitClauseWithVarList(Clause);
+}
+
+void OpenACCClauseProfiler::VisitDeviceResidentClause(
+    const OpenACCDeviceResidentClause &Clause) {
+  VisitClauseWithVarList(Clause);
+}
+
 void OpenACCClauseProfiler::VisitCopyInClause(
     const OpenACCCopyInClause &Clause) {
   VisitClauseWithVarList(Clause);

diff  --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp
index fd1eaab9621dd..063e718454d46 100644
--- a/clang/lib/AST/TextNodeDumper.cpp
+++ b/clang/lib/AST/TextNodeDumper.cpp
@@ -417,9 +417,11 @@ void TextNodeDumper::Visit(const OpenACCClause *C) {
     case OpenACCClauseKind::Device:
     case OpenACCClauseKind::DeviceNum:
     case OpenACCClauseKind::DefaultAsync:
+    case OpenACCClauseKind::DeviceResident:
     case OpenACCClauseKind::DevicePtr:
     case OpenACCClauseKind::Finalize:
     case OpenACCClauseKind::FirstPrivate:
+    case OpenACCClauseKind::Link:
     case OpenACCClauseKind::NoCreate:
     case OpenACCClauseKind::NumGangs:
     case OpenACCClauseKind::NumWorkers:
@@ -3061,6 +3063,17 @@ void TextNodeDumper::VisitOpenACCAtomicConstruct(
   OS << ' ' << S->getAtomicKind();
 }
 
+void TextNodeDumper::VisitOpenACCDeclareDecl(const OpenACCDeclareDecl *D) {
+  OS << " " << D->getDirectiveKind();
+
+  for (const OpenACCClause *C : D->clauses())
+    AddChild([=] {
+      Visit(C);
+      for (const Stmt *S : C->children())
+        AddChild([=] { Visit(S); });
+    });
+}
+
 void TextNodeDumper::VisitEmbedExpr(const EmbedExpr *S) {
   AddChild("begin", [=] { OS << S->getStartingElementPos(); });
   AddChild("number of elements", [=] { OS << S->getDataElementCount(); });

diff  --git a/clang/lib/CodeGen/CGDecl.cpp b/clang/lib/CodeGen/CGDecl.cpp
index 668282a6ab1a8..9cd5885aaae51 100644
--- a/clang/lib/CodeGen/CGDecl.cpp
+++ b/clang/lib/CodeGen/CGDecl.cpp
@@ -27,6 +27,7 @@
 #include "clang/AST/CharUnits.h"
 #include "clang/AST/Decl.h"
 #include "clang/AST/DeclObjC.h"
+#include "clang/AST/DeclOpenACC.h"
 #include "clang/AST/DeclOpenMP.h"
 #include "clang/Basic/CodeGenOptions.h"
 #include "clang/Basic/TargetInfo.h"
@@ -177,6 +178,9 @@ void CodeGenFunction::EmitDecl(const Decl &D) {
   case Decl::OMPDeclareMapper:
     return CGM.EmitOMPDeclareMapper(cast<OMPDeclareMapperDecl>(&D), this);
 
+  case Decl::OpenACCDeclare:
+    return CGM.EmitOpenACCDeclare(cast<OpenACCDeclareDecl>(&D), this);
+
   case Decl::Typedef:      // typedef int X;
   case Decl::TypeAlias: {  // using X = int; [C++0x]
     QualType Ty = cast<TypedefNameDecl>(D).getUnderlyingType();
@@ -2843,6 +2847,11 @@ void CodeGenModule::EmitOMPDeclareMapper(const OMPDeclareMapperDecl *D,
   getOpenMPRuntime().emitUserDefinedMapper(D, CGF);
 }
 
+void CodeGenModule::EmitOpenACCDeclare(const OpenACCDeclareDecl *D,
+                                       CodeGenFunction *CGF) {
+  // This is a no-op, we cna just ignore these declarations.
+}
+
 void CodeGenModule::EmitOMPRequiresDecl(const OMPRequiresDecl *D) {
   getOpenMPRuntime().processRequiresDirective(D);
 }

diff  --git a/clang/lib/CodeGen/CodeGenModule.h b/clang/lib/CodeGen/CodeGenModule.h
index 4a269f622ece4..5726d7939b3c3 100644
--- a/clang/lib/CodeGen/CodeGenModule.h
+++ b/clang/lib/CodeGen/CodeGenModule.h
@@ -1569,6 +1569,9 @@ class CodeGenModule : public CodeGenTypeCache {
   void EmitOMPDeclareMapper(const OMPDeclareMapperDecl *D,
                             CodeGenFunction *CGF = nullptr);
 
+  // Emit code for the OpenACC Declare declaration.
+  void EmitOpenACCDeclare(const OpenACCDeclareDecl *D, CodeGenFunction *CGF);
+
   /// Emit a code for requires directive.
   /// \param D Requires declaration
   void EmitOMPRequiresDecl(const OMPRequiresDecl *D);

diff  --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp
index d036971d2fc31..5eefd3ffb092e 100644
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@@ -981,7 +981,7 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
     case OpenACCClauseKind::PresentOrCopyIn: {
       bool IsReadOnly = tryParseAndConsumeSpecialTokenKind(
           *this, OpenACCSpecialTokenKind::ReadOnly, ClauseKind);
-      ParsedClause.setVarListDetails(ParseOpenACCVarList(ClauseKind),
+      ParsedClause.setVarListDetails(ParseOpenACCVarList(DirKind, ClauseKind),
                                      IsReadOnly,
                                      /*IsZero=*/false);
       break;
@@ -994,7 +994,7 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
     case OpenACCClauseKind::PresentOrCopyOut: {
       bool IsZero = tryParseAndConsumeSpecialTokenKind(
           *this, OpenACCSpecialTokenKind::Zero, ClauseKind);
-      ParsedClause.setVarListDetails(ParseOpenACCVarList(ClauseKind),
+      ParsedClause.setVarListDetails(ParseOpenACCVarList(DirKind, ClauseKind),
                                      /*IsReadOnly=*/false, IsZero);
       break;
     }
@@ -1002,7 +1002,8 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
       // If we're missing a clause-kind (or it is invalid), see if we can parse
       // the var-list anyway.
       OpenACCReductionOperator Op = ParseReductionOperator(*this);
-      ParsedClause.setReductionDetails(Op, ParseOpenACCVarList(ClauseKind));
+      ParsedClause.setReductionDetails(
+          Op, ParseOpenACCVarList(DirKind, ClauseKind));
       break;
     }
     case OpenACCClauseKind::Self:
@@ -1013,21 +1014,13 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
       [[fallthrough]];
     case OpenACCClauseKind::Device:
     case OpenACCClauseKind::Host:
-      ParsedClause.setVarListDetails(ParseOpenACCVarList(ClauseKind),
-                                     /*IsReadOnly=*/false, /*IsZero=*/false);
-      break;
     case OpenACCClauseKind::DeviceResident:
     case OpenACCClauseKind::Link:
-      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;
     case OpenACCClauseKind::Copy:
     case OpenACCClauseKind::PCopy:
     case OpenACCClauseKind::PresentOrCopy:
@@ -1035,7 +1028,7 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
     case OpenACCClauseKind::NoCreate:
     case OpenACCClauseKind::Present:
     case OpenACCClauseKind::Private:
-      ParsedClause.setVarListDetails(ParseOpenACCVarList(ClauseKind),
+      ParsedClause.setVarListDetails(ParseOpenACCVarList(DirKind, ClauseKind),
                                      /*IsReadOnly=*/false, /*IsZero=*/false);
       break;
     case OpenACCClauseKind::Collapse: {
@@ -1362,7 +1355,8 @@ ExprResult Parser::ParseOpenACCBindClauseArgument() {
 /// - an array element
 /// - a member of a composite variable
 /// - a common block name between slashes (fortran only)
-Parser::OpenACCVarParseResult Parser::ParseOpenACCVar(OpenACCClauseKind CK) {
+Parser::OpenACCVarParseResult Parser::ParseOpenACCVar(OpenACCDirectiveKind DK,
+                                                      OpenACCClauseKind CK) {
   OpenACCArraySectionRAII ArraySections(*this);
 
   ExprResult Res = ParseAssignmentExpression();
@@ -1373,15 +1367,16 @@ Parser::OpenACCVarParseResult Parser::ParseOpenACCVar(OpenACCClauseKind CK) {
   if (!Res.isUsable())
     return {Res, OpenACCParseCanContinue::Can};
 
-  Res = getActions().OpenACC().ActOnVar(CK, Res.get());
+  Res = getActions().OpenACC().ActOnVar(DK, CK, Res.get());
 
   return {Res, OpenACCParseCanContinue::Can};
 }
 
-llvm::SmallVector<Expr *> Parser::ParseOpenACCVarList(OpenACCClauseKind CK) {
+llvm::SmallVector<Expr *> Parser::ParseOpenACCVarList(OpenACCDirectiveKind DK,
+                                                      OpenACCClauseKind CK) {
   llvm::SmallVector<Expr *> Vars;
 
-  auto [Res, CanContinue] = ParseOpenACCVar(CK);
+  auto [Res, CanContinue] = ParseOpenACCVar(DK, CK);
   if (Res.isUsable()) {
     Vars.push_back(Res.get());
   } else if (CanContinue == OpenACCParseCanContinue::Cannot) {
@@ -1392,7 +1387,7 @@ llvm::SmallVector<Expr *> Parser::ParseOpenACCVarList(OpenACCClauseKind CK) {
   while (!getCurToken().isOneOf(tok::r_paren, tok::annot_pragma_openacc_end)) {
     ExpectAndConsume(tok::comma);
 
-    auto [Res, CanContinue] = ParseOpenACCVar(CK);
+    auto [Res, CanContinue] = ParseOpenACCVar(DK, CK);
 
     if (Res.isUsable()) {
       Vars.push_back(Res.get());
@@ -1426,7 +1421,7 @@ void Parser::ParseOpenACCCacheVarList() {
 
   // ParseOpenACCVarList should leave us before a r-paren, so no need to skip
   // anything here.
-  ParseOpenACCVarList(OpenACCClauseKind::Invalid);
+  ParseOpenACCVarList(OpenACCDirectiveKind::Cache, OpenACCClauseKind::Invalid);
 }
 
 Parser::OpenACCDirectiveParseInfo
@@ -1523,8 +1518,9 @@ Parser::DeclGroupPtrTy Parser::ParseOpenACCDirectiveDecl() {
                                                      DirInfo.StartLoc))
     return nullptr;
 
-  // TODO OpenACC: Do whatever decl parsing is required here.
-  return DeclGroupPtrTy::make(getActions().OpenACC().ActOnEndDeclDirective());
+  return DeclGroupPtrTy::make(getActions().OpenACC().ActOnEndDeclDirective(
+      DirInfo.DirKind, DirInfo.StartLoc, DirInfo.DirLoc, DirInfo.EndLoc,
+      DirInfo.Clauses));
 }
 
 // Parse OpenACC Directive on a Statement.

diff  --git a/clang/lib/Sema/SemaBase.cpp b/clang/lib/Sema/SemaBase.cpp
index 5c24f21b469b0..85c4a0ab40fed 100644
--- a/clang/lib/Sema/SemaBase.cpp
+++ b/clang/lib/Sema/SemaBase.cpp
@@ -9,6 +9,7 @@ SemaBase::SemaBase(Sema &S) : SemaRef(S) {}
 ASTContext &SemaBase::getASTContext() const { return SemaRef.Context; }
 DiagnosticsEngine &SemaBase::getDiagnostics() const { return SemaRef.Diags; }
 const LangOptions &SemaBase::getLangOpts() const { return SemaRef.LangOpts; }
+DeclContext *SemaBase::getCurContext() const { return SemaRef.CurContext; }
 
 SemaBase::ImmediateDiagBuilder::~ImmediateDiagBuilder() {
   // If we aren't active, there is nothing to do.

diff  --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 2d2f8ddf4652b..59583e73952cf 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -11,11 +11,12 @@
 ///
 //===----------------------------------------------------------------------===//
 
+#include "clang/Sema/SemaOpenACC.h"
+#include "clang/AST/DeclOpenACC.h"
 #include "clang/AST/StmtOpenACC.h"
 #include "clang/Basic/DiagnosticSema.h"
 #include "clang/Basic/OpenACCKinds.h"
 #include "clang/Sema/Sema.h"
-#include "clang/Sema/SemaOpenACC.h"
 #include "llvm/ADT/StringExtras.h"
 #include "llvm/Support/Casting.h"
 
@@ -334,6 +335,7 @@ void SemaOpenACC::ActOnConstruct(OpenACCDirectiveKind K,
   case OpenACCDirectiveKind::Set:
   case OpenACCDirectiveKind::Update:
   case OpenACCDirectiveKind::Atomic:
+  case OpenACCDirectiveKind::Declare:
     // Nothing to do here, there is no real legalization that needs to happen
     // here as these constructs do not take any arguments.
     break;
@@ -478,15 +480,20 @@ bool SemaOpenACC::CheckVarIsPointerType(OpenACCClauseKind ClauseKind,
   return false;
 }
 
-ExprResult SemaOpenACC::ActOnVar(OpenACCClauseKind CK, Expr *VarExpr) {
+ExprResult SemaOpenACC::ActOnVar(OpenACCDirectiveKind DK, 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 &&
+  // OpenACC3.3 2.13:
+  // A 'var' in a 'declare' directive must be a variable or array name.
+  if ((CK == OpenACCClauseKind::UseDevice ||
+       DK == OpenACCDirectiveKind::Declare) &&
       isa<ArraySectionExpr, ArraySubscriptExpr>(CurVarExpr)) {
-    Diag(VarExpr->getExprLoc(), diag::err_acc_not_a_var_ref_use_device);
+    Diag(VarExpr->getExprLoc(), diag::err_acc_not_a_var_ref_use_device_declare)
+        << (DK == OpenACCDirectiveKind::Declare);
     return ExprError();
   }
 
@@ -510,20 +517,30 @@ 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.
-  // If CK is a 'use_device', this also isn't valid, as it isn' the name of a
-  // variable or array.
+  // If CK is a 'use_device', this also isn't valid, as it isn't the name of a
+  // variable or array, if not done as a member expr.
   // 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()))
+  if (const auto *ME = dyn_cast<MemberExpr>(CurVarExpr)) {
+    if (isa<FieldDecl>(ME->getMemberDecl()->getCanonicalDecl())) {
+      if (DK == OpenACCDirectiveKind::Declare ||
+          CK == OpenACCClauseKind::Reduction ||
+          CK == OpenACCClauseKind::UseDevice) {
+
+        // We can allow 'member expr' if the 'this' is implicit in the case of
+        // declare, reduction, and use_device.
+        const auto *This = dyn_cast<CXXThisExpr>(ME->getBase());
+        if (This && This->isImplicit())
+          return VarExpr;
+      } else {
         return VarExpr;
+      }
     }
   }
 
-  // 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))
+  // Referring to 'this' is ok for the most part, but for 'use_device'/'declare'
+  // doesn't fall into 'variable or array name'
+  if (CK != OpenACCClauseKind::UseDevice &&
+      DK != OpenACCDirectiveKind::Declare && isa<CXXThisExpr>(CurVarExpr))
     return VarExpr;
 
   // Nothing really we can do here, as these are dependent.  So just return they
@@ -538,8 +555,12 @@ ExprResult SemaOpenACC::ActOnVar(OpenACCClauseKind CK, Expr *VarExpr) {
   if (isa<RecoveryExpr>(CurVarExpr))
     return ExprError();
 
-  if (CK == OpenACCClauseKind::UseDevice)
-    Diag(VarExpr->getExprLoc(), diag::err_acc_not_a_var_ref_use_device);
+  if (DK == OpenACCDirectiveKind::Declare)
+    Diag(VarExpr->getExprLoc(), diag::err_acc_not_a_var_ref_use_device_declare)
+        << /*declare*/ 1;
+  else if (CK == OpenACCClauseKind::UseDevice)
+    Diag(VarExpr->getExprLoc(), diag::err_acc_not_a_var_ref_use_device_declare)
+        << /*use_device*/ 0;
   else
     Diag(VarExpr->getExprLoc(), diag::err_acc_not_a_var_ref)
         << (CK != OpenACCClauseKind::Reduction);
@@ -1413,6 +1434,12 @@ std::string GetListOfClauses(llvm::ArrayRef<OpenACCClauseKind> Clauses) {
 bool SemaOpenACC::ActOnStartStmtDirective(
     OpenACCDirectiveKind K, SourceLocation StartLoc,
     ArrayRef<const OpenACCClause *> Clauses) {
+
+  // Declaration directives an appear in a statement location, so call into that
+  // function here.
+  if (K == OpenACCDirectiveKind::Declare || K == OpenACCDirectiveKind::Routine)
+    return ActOnStartDeclDirective(K, StartLoc);
+
   SemaRef.DiscardCleanupsInEvaluationContext();
   SemaRef.PopExpressionEvaluationContext();
 
@@ -1597,6 +1624,14 @@ StmtResult SemaOpenACC::ActOnEndStmtDirective(
         getASTContext(), StartLoc, DirLoc, AtomicKind, EndLoc,
         AssocStmt.isUsable() ? AssocStmt.get() : nullptr);
   }
+  case OpenACCDirectiveKind::Declare: {
+    // Declare is a declaration directive, but can be used here as long as we
+    // wrap it in a DeclStmt.  So make sure we do that here.
+
+    DeclGroupRef DR =
+        ActOnEndDeclDirective(K, StartLoc, DirLoc, EndLoc, Clauses);
+    return SemaRef.ActOnDeclStmt(DeclGroupPtrTy::make(DR), StartLoc, EndLoc);
+  }
   }
   llvm_unreachable("Unhandled case in directive handling?");
 }
@@ -1683,7 +1718,33 @@ bool SemaOpenACC::ActOnStartDeclDirective(OpenACCDirectiveKind K,
   return diagnoseConstructAppertainment(*this, K, StartLoc, /*IsStmt=*/false);
 }
 
-DeclGroupRef SemaOpenACC::ActOnEndDeclDirective() { return DeclGroupRef{}; }
+DeclGroupRef SemaOpenACC::ActOnEndDeclDirective(
+    OpenACCDirectiveKind K, SourceLocation StartLoc, SourceLocation DirLoc,
+    SourceLocation EndLoc, ArrayRef<OpenACCClause *> Clauses) {
+  switch (K) {
+  default:
+  case OpenACCDirectiveKind::Invalid:
+    return DeclGroupRef{};
+  case OpenACCDirectiveKind::Declare: {
+    // OpenACC3.3 2.13: At least one clause must appear on a declare directive.
+    if (Clauses.empty()) {
+      Diag(EndLoc, diag::err_acc_declare_required_clauses);
+      // No reason to add this to the AST, as we would just end up trying to
+      // instantiate this, which would double-diagnose here, which we wouldn't
+      // want to do.
+      return DeclGroupRef{};
+    }
+
+    auto *DeclareDecl = OpenACCDeclareDecl::Create(
+        getASTContext(), getCurContext(), StartLoc, DirLoc, EndLoc, Clauses);
+    DeclareDecl->setAccess(AS_public);
+    getCurContext()->addDecl(DeclareDecl);
+    return DeclGroupRef{DeclareDecl};
+  }
+  }
+
+  llvm_unreachable("unhandled case in directive handling?");
+}
 
 ExprResult
 SemaOpenACC::BuildOpenACCAsteriskSizeExpr(SourceLocation AsteriskLoc) {

diff  --git a/clang/lib/Sema/SemaOpenACCClause.cpp b/clang/lib/Sema/SemaOpenACCClause.cpp
index 1e74f126c31ce..0805779d74bdc 100644
--- a/clang/lib/Sema/SemaOpenACCClause.cpp
+++ b/clang/lib/Sema/SemaOpenACCClause.cpp
@@ -10,8 +10,9 @@
 ///
 //===----------------------------------------------------------------------===//
 
-#include "clang/AST/OpenACCClause.h"
 #include "clang/AST/DeclCXX.h"
+#include "clang/AST/ExprCXX.h"
+#include "clang/AST/OpenACCClause.h"
 #include "clang/Basic/DiagnosticSema.h"
 #include "clang/Basic/OpenACCKinds.h"
 #include "clang/Sema/SemaOpenACC.h"
@@ -192,6 +193,7 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
     case OpenACCDirectiveKind::Kernels:
     case OpenACCDirectiveKind::Data:
     case OpenACCDirectiveKind::EnterData:
+    case OpenACCDirectiveKind::Declare:
     case OpenACCDirectiveKind::ParallelLoop:
     case OpenACCDirectiveKind::SerialLoop:
     case OpenACCDirectiveKind::KernelsLoop:
@@ -424,6 +426,22 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
       return false;
     }
   }
+  case OpenACCClauseKind::Link: {
+    switch (DirectiveKind) {
+    case OpenACCDirectiveKind::Declare:
+      return true;
+    default:
+      return false;
+    }
+  }
+  case OpenACCClauseKind::DeviceResident: {
+    switch (DirectiveKind) {
+    case OpenACCDirectiveKind::Declare:
+      return true;
+    default:
+      return false;
+    }
+  }
 
   case OpenACCClauseKind::UseDevice: {
     switch (DirectiveKind) {
@@ -588,8 +606,17 @@ bool checkValidAfterDeviceType(
 // with the one being currently implemented/only updated after the entire
 // construct has been implemented.
 bool isDirectiveKindImplemented(OpenACCDirectiveKind DK) {
-  return DK != OpenACCDirectiveKind::Declare &&
-         DK != OpenACCDirectiveKind::Routine;
+  return DK != OpenACCDirectiveKind::Routine;
+}
+
+// GCC looks through linkage specs, but not the other transparent declaration
+// contexts for 'declare' restrictions, so this helper function helps get us
+// through that.
+const DeclContext *removeLinkageSpecDC(const DeclContext *DC) {
+  while (isa<LinkageSpecDecl>(DC))
+    DC = DC->getParent();
+
+  return DC;
 }
 
 class SemaOpenACCClauseVisitor {
@@ -1006,16 +1033,15 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitNoCreateClause(
 
 OpenACCClause *SemaOpenACCClauseVisitor::VisitPresentClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
-  // Restrictions only properly implemented on 'compute'/'combined'/'data'
-  // constructs, and 'compute'/'combined'/'data' constructs are the only
-  // construct that can do anything with this yet, so skip/treat as
-  // unimplemented in this case.
-  if (!isDirectiveKindImplemented(Clause.getDirectiveKind()))
-    return isNotImplemented();
   // ActOnVar ensured that everything is a valid variable reference, so there
   // really isn't anything to do here. GCC does some duplicate-finding, though
   // it isn't apparent in the standard where this is justified.
 
+  // 'declare' has some restrictions that need to be enforced separately, so
+  // check it here.
+  if (SemaRef.CheckDeclareClause(Clause))
+    return nullptr;
+
   return OpenACCPresentClause::Create(Ctx, Clause.getBeginLoc(),
                                       Clause.getLParenLoc(),
                                       Clause.getVarList(), Clause.getEndLoc());
@@ -1045,33 +1071,58 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitDeviceClause(
 
 OpenACCClause *SemaOpenACCClauseVisitor::VisitCopyClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
-  // Restrictions only properly implemented on 'compute'/'combined'/'data'
-  // constructs, and 'compute'/'combined'/'data' constructs are the only
-  // construct that can do anything with this yet, so skip/treat as
-  // unimplemented in this case.
-  if (!isDirectiveKindImplemented(Clause.getDirectiveKind()))
-    return isNotImplemented();
   // ActOnVar ensured that everything is a valid variable reference, so there
   // really isn't anything to do here. GCC does some duplicate-finding, though
   // it isn't apparent in the standard where this is justified.
 
+  // 'declare' has some restrictions that need to be enforced separately, so
+  // check it here.
+  if (SemaRef.CheckDeclareClause(Clause))
+    return nullptr;
+
   return OpenACCCopyClause::Create(
       Ctx, Clause.getClauseKind(), Clause.getBeginLoc(), Clause.getLParenLoc(),
       Clause.getVarList(), Clause.getEndLoc());
 }
 
+OpenACCClause *SemaOpenACCClauseVisitor::VisitLinkClause(
+    SemaOpenACC::OpenACCParsedClause &Clause) {
+  // 'declare' has some restrictions that need to be enforced separately, so
+  // check it here.
+  if (SemaRef.CheckDeclareClause(Clause))
+    return nullptr;
+
+  Clause.setVarListDetails(SemaRef.CheckLinkClauseVarList(Clause.getVarList()),
+                           /*IsReadOnly=*/false, /*IsZero=*/false);
+
+  return OpenACCLinkClause::Create(Ctx, Clause.getBeginLoc(),
+                                   Clause.getLParenLoc(), Clause.getVarList(),
+                                   Clause.getEndLoc());
+}
+
+OpenACCClause *SemaOpenACCClauseVisitor::VisitDeviceResidentClause(
+    SemaOpenACC::OpenACCParsedClause &Clause) {
+  // 'declare' has some restrictions that need to be enforced separately, so
+  // check it here.
+  if (SemaRef.CheckDeclareClause(Clause))
+    return nullptr;
+
+  return OpenACCDeviceResidentClause::Create(
+      Ctx, Clause.getBeginLoc(), Clause.getLParenLoc(), Clause.getVarList(),
+      Clause.getEndLoc());
+}
+
 OpenACCClause *SemaOpenACCClauseVisitor::VisitCopyInClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
-  // Restrictions only properly implemented on 'compute'/'combined'/'data'
-  // constructs, and 'compute'/'combined'/'data' constructs are the only
-  // construct that can do anything with this yet, so skip/treat as
-  // unimplemented in this case.
-  if (!isDirectiveKindImplemented(Clause.getDirectiveKind()))
-    return isNotImplemented();
   // ActOnVar ensured that everything is a valid variable reference, so there
   // really isn't anything to do here. GCC does some duplicate-finding, though
   // it isn't apparent in the standard where this is justified.
 
+  // 'declare' has some restrictions that need to be enforced separately, so
+  // check it here.
+  if (SemaRef.CheckDeclareClause(Clause))
+    return nullptr;
+
   return OpenACCCopyInClause::Create(
       Ctx, Clause.getClauseKind(), Clause.getBeginLoc(), Clause.getLParenLoc(),
       Clause.isReadOnly(), Clause.getVarList(), Clause.getEndLoc());
@@ -1079,16 +1130,15 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitCopyInClause(
 
 OpenACCClause *SemaOpenACCClauseVisitor::VisitCopyOutClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
-  // Restrictions only properly implemented on 'compute'/'combined'/'data'
-  // constructs, and 'compute'/'combined'/'data' constructs are the only
-  // construct that can do anything with this yet, so skip/treat as
-  // unimplemented in this case.
-  if (!isDirectiveKindImplemented(Clause.getDirectiveKind()))
-    return isNotImplemented();
   // ActOnVar ensured that everything is a valid variable reference, so there
   // really isn't anything to do here. GCC does some duplicate-finding, though
   // it isn't apparent in the standard where this is justified.
 
+  // 'declare' has some restrictions that need to be enforced separately, so
+  // check it here.
+  if (SemaRef.CheckDeclareClause(Clause))
+    return nullptr;
+
   return OpenACCCopyOutClause::Create(
       Ctx, Clause.getClauseKind(), Clause.getBeginLoc(), Clause.getLParenLoc(),
       Clause.isZero(), Clause.getVarList(), Clause.getEndLoc());
@@ -1100,6 +1150,11 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitCreateClause(
   // really isn't anything to do here. GCC does some duplicate-finding, though
   // it isn't apparent in the standard where this is justified.
 
+  // 'declare' has some restrictions that need to be enforced separately, so
+  // check it here.
+  if (SemaRef.CheckDeclareClause(Clause))
+    return nullptr;
+
   return OpenACCCreateClause::Create(
       Ctx, Clause.getClauseKind(), Clause.getBeginLoc(), Clause.getLParenLoc(),
       Clause.isZero(), Clause.getVarList(), Clause.getEndLoc());
@@ -1156,13 +1211,6 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitUseDeviceClause(
 
 OpenACCClause *SemaOpenACCClauseVisitor::VisitDevicePtrClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
-  // Restrictions only properly implemented on 'compute'/'combined'/'data'
-  // constructs, and 'compute'/'combined'/'data' constructs are the only
-  // construct that can do anything with this yet, so skip/treat as
-  // unimplemented in this case.
-  if (!isDirectiveKindImplemented(Clause.getDirectiveKind()))
-    return isNotImplemented();
-
   // ActOnVar ensured that everything is a valid variable reference, but we
   // still have to make sure it is a pointer type.
   llvm::SmallVector<Expr *> VarList{Clause.getVarList()};
@@ -1172,6 +1220,11 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitDevicePtrClause(
   Clause.setVarListDetails(VarList,
                            /*IsReadOnly=*/false, /*IsZero=*/false);
 
+  // 'declare' has some restrictions that need to be enforced separately, so
+  // check it here.
+  if (SemaRef.CheckDeclareClause(Clause))
+    return nullptr;
+
   return OpenACCDevicePtrClause::Create(
       Ctx, Clause.getBeginLoc(), Clause.getLParenLoc(), Clause.getVarList(),
       Clause.getEndLoc());
@@ -2243,3 +2296,136 @@ OpenACCClause *SemaOpenACC::CheckReductionClause(
       getASTContext(), BeginLoc, LParenLoc, ReductionOp, Vars, EndLoc);
   return Ret;
 }
+
+llvm::SmallVector<Expr *>
+SemaOpenACC::CheckLinkClauseVarList(ArrayRef<Expr *> VarExprs) {
+  const DeclContext *DC = removeLinkageSpecDC(getCurContext());
+
+  // Link has no special restrictions on its var list unless it is not at NS/TU
+  // scope.
+  if (isa<NamespaceDecl, TranslationUnitDecl>(DC))
+    return llvm::SmallVector<Expr *>(VarExprs);
+
+  llvm::SmallVector<Expr *> NewVarList;
+
+  for (Expr *VarExpr : VarExprs) {
+    if (isa<DependentScopeDeclRefExpr, CXXDependentScopeMemberExpr>(VarExpr)) {
+      NewVarList.push_back(VarExpr);
+      continue;
+    }
+
+    // Field decls can't be global, nor extern, and declare can't refer to
+    // non-static fields in class-scope, so this always fails the scope check.
+    // BUT for now we add this so it gets diagnosed by the general 'declare'
+    // rules.
+    if (isa<MemberExpr>(VarExpr)) {
+      NewVarList.push_back(VarExpr);
+      continue;
+    }
+
+    const auto *DRE = cast<DeclRefExpr>(VarExpr);
+    const VarDecl *Var = dyn_cast<VarDecl>(DRE->getDecl());
+
+    if (!Var || !Var->hasExternalStorage())
+      Diag(VarExpr->getBeginLoc(), diag::err_acc_link_not_extern);
+    else
+      NewVarList.push_back(VarExpr);
+  }
+
+  return NewVarList;
+}
+bool SemaOpenACC::CheckDeclareClause(SemaOpenACC::OpenACCParsedClause &Clause) {
+
+  if (Clause.getDirectiveKind() != OpenACCDirectiveKind::Declare)
+    return false;
+
+  const DeclContext *DC = removeLinkageSpecDC(getCurContext());
+
+  // Whether this is 'create', 'copyin', 'deviceptr', 'device_resident', or
+  // 'link', which have 2 special rules.
+  bool IsSpecialClause =
+      Clause.getClauseKind() == OpenACCClauseKind::Create ||
+      Clause.getClauseKind() == OpenACCClauseKind::CopyIn ||
+      Clause.getClauseKind() == OpenACCClauseKind::DevicePtr ||
+      Clause.getClauseKind() == OpenACCClauseKind::DeviceResident ||
+      Clause.getClauseKind() == OpenACCClauseKind::Link;
+
+  // OpenACC 3.3 2.13:
+  // In C or C++ global or namespace scope, only 'create',
+  // 'copyin', 'deviceptr', 'device_resident', or 'link' clauses are
+  // allowed.
+  if (!IsSpecialClause && isa<NamespaceDecl, TranslationUnitDecl>(DC)) {
+    return Diag(Clause.getBeginLoc(), diag::err_acc_declare_clause_at_global)
+           << Clause.getClauseKind();
+  }
+
+  llvm::SmallVector<Expr *> FilteredVarList;
+  const DeclaratorDecl *CurDecl = nullptr;
+  for (Expr *VarExpr : Clause.getVarList()) {
+    if (isa<DependentScopeDeclRefExpr, CXXDependentScopeMemberExpr>(VarExpr)) {
+      // There isn't really anything we can do here, so we add them anyway and
+      // we can check them again when we instantiate this.
+    } else if (const auto *MemExpr = dyn_cast<MemberExpr>(VarExpr)) {
+      FieldDecl *FD =
+          cast<FieldDecl>(MemExpr->getMemberDecl()->getCanonicalDecl());
+      CurDecl = FD;
+
+      if (removeLinkageSpecDC(
+              FD->getLexicalDeclContext()->getPrimaryContext()) != DC) {
+        Diag(MemExpr->getBeginLoc(), diag::err_acc_declare_same_scope)
+            << Clause.getClauseKind();
+        continue;
+      }
+    } else {
+      const auto *DRE = cast<DeclRefExpr>(VarExpr);
+      const VarDecl *Var = dyn_cast<VarDecl>(DRE->getDecl());
+      if (Var)
+        CurDecl = Var->getCanonicalDecl();
+
+      // OpenACC3.3 2.13:
+      // A 'declare' directive must be in the same scope as the declaration of
+      // any var that appears in the clauses of the directive or any scope
+      // within a C/C++ function.
+      // We can't really check 'scope' here, so we check declaration context,
+      // which is a reasonable approximation, but misses scopes inside of
+      // functions.
+      if (removeLinkageSpecDC(Var->getCanonicalDecl()
+                                  ->getLexicalDeclContext()
+                                  ->getPrimaryContext()) != DC) {
+        Diag(VarExpr->getBeginLoc(), diag::err_acc_declare_same_scope)
+            << Clause.getClauseKind();
+        continue;
+      }
+      // OpenACC3.3 2.13:
+      // C and C++ extern variables may only appear in 'create',
+      // 'copyin', 'deviceptr', 'device_resident', or 'link' clauses on a
+      // 'declare' directive.
+      if (!IsSpecialClause && Var && Var->hasExternalStorage()) {
+        Diag(VarExpr->getBeginLoc(), diag::err_acc_declare_extern)
+            << Clause.getClauseKind();
+        continue;
+      }
+
+      // OpenACC3.3 2.13:
+      // A var may appear at most once in all the clauses of declare
+      // directives for a function, subroutine, program, or module.
+
+      if (CurDecl) {
+        auto Itr = DeclareVarReferences.find(CurDecl);
+        if (Itr != DeclareVarReferences.end()) {
+          Diag(VarExpr->getBeginLoc(), diag::err_acc_multiple_references)
+              << Clause.getClauseKind();
+          Diag(Itr->second, diag::note_acc_previous_reference);
+          continue;
+        } else {
+          DeclareVarReferences[CurDecl] = VarExpr->getBeginLoc();
+        }
+      }
+    }
+    FilteredVarList.push_back(VarExpr);
+  }
+
+  Clause.setVarListDetails(FilteredVarList, Clause.isReadOnly(),
+                           Clause.isZero());
+  return false;
+}

diff  --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
index 1cdf80898bfca..dd894df851488 100644
--- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
+++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
@@ -999,6 +999,234 @@ TemplateDeclInstantiator::VisitNamespaceDecl(NamespaceDecl *D) {
   llvm_unreachable("Namespaces cannot be instantiated");
 }
 
+namespace {
+class OpenACCDeclClauseInstantiator final
+    : public OpenACCClauseVisitor<OpenACCDeclClauseInstantiator> {
+  Sema &SemaRef;
+  const MultiLevelTemplateArgumentList &MLTAL;
+  SemaOpenACC::OpenACCParsedClause &ParsedClause;
+  OpenACCClause *NewClause = nullptr;
+
+public:
+  OpenACCDeclClauseInstantiator(Sema &S,
+                                const MultiLevelTemplateArgumentList &MLTAL,
+                                SemaOpenACC::OpenACCParsedClause &ParsedClause)
+      : SemaRef(S), MLTAL(MLTAL), ParsedClause(ParsedClause) {}
+
+  OpenACCClause *CreatedClause() { return NewClause; }
+#define VISIT_CLAUSE(CLAUSE_NAME)                                              \
+  void Visit##CLAUSE_NAME##Clause(const OpenACC##CLAUSE_NAME##Clause &Clause);
+#include "clang/Basic/OpenACCClauses.def"
+
+  llvm::SmallVector<Expr *> VisitVarList(ArrayRef<Expr *> VarList) {
+    llvm::SmallVector<Expr *> InstantiatedVarList;
+    for (Expr *CurVar : VarList) {
+      ExprResult Res = SemaRef.SubstExpr(CurVar, MLTAL);
+
+      if (!Res.isUsable())
+        continue;
+
+      Res = SemaRef.OpenACC().ActOnVar(ParsedClause.getDirectiveKind(),
+                                       ParsedClause.getClauseKind(), Res.get());
+
+      if (Res.isUsable())
+        InstantiatedVarList.push_back(Res.get());
+    }
+    return InstantiatedVarList;
+  }
+};
+
+#define CLAUSE_NOT_ON_DECLS(CLAUSE_NAME)                                       \
+  void OpenACCDeclClauseInstantiator::Visit##CLAUSE_NAME##Clause(              \
+      const OpenACC##CLAUSE_NAME##Clause &) {                                  \
+    llvm_unreachable("Clause type invalid on declaration construct, or "       \
+                     "instantiation not implemented");                         \
+  }
+
+CLAUSE_NOT_ON_DECLS(Auto)
+CLAUSE_NOT_ON_DECLS(Async)
+CLAUSE_NOT_ON_DECLS(Attach)
+CLAUSE_NOT_ON_DECLS(Collapse)
+CLAUSE_NOT_ON_DECLS(Default)
+CLAUSE_NOT_ON_DECLS(DefaultAsync)
+CLAUSE_NOT_ON_DECLS(Delete)
+CLAUSE_NOT_ON_DECLS(Detach)
+CLAUSE_NOT_ON_DECLS(Device)
+CLAUSE_NOT_ON_DECLS(DeviceNum)
+CLAUSE_NOT_ON_DECLS(DeviceType)
+CLAUSE_NOT_ON_DECLS(Finalize)
+CLAUSE_NOT_ON_DECLS(FirstPrivate)
+CLAUSE_NOT_ON_DECLS(Gang)
+CLAUSE_NOT_ON_DECLS(Host)
+CLAUSE_NOT_ON_DECLS(If)
+CLAUSE_NOT_ON_DECLS(IfPresent)
+CLAUSE_NOT_ON_DECLS(Independent)
+CLAUSE_NOT_ON_DECLS(NoCreate)
+CLAUSE_NOT_ON_DECLS(NumGangs)
+CLAUSE_NOT_ON_DECLS(NumWorkers)
+CLAUSE_NOT_ON_DECLS(Private)
+CLAUSE_NOT_ON_DECLS(Reduction)
+CLAUSE_NOT_ON_DECLS(Self)
+CLAUSE_NOT_ON_DECLS(Seq)
+CLAUSE_NOT_ON_DECLS(Tile)
+CLAUSE_NOT_ON_DECLS(UseDevice)
+CLAUSE_NOT_ON_DECLS(Vector)
+CLAUSE_NOT_ON_DECLS(VectorLength)
+CLAUSE_NOT_ON_DECLS(Wait)
+CLAUSE_NOT_ON_DECLS(Worker)
+#undef CLAUSE_NOT_ON_DECLS
+
+void OpenACCDeclClauseInstantiator::VisitCopyClause(
+    const OpenACCCopyClause &C) {
+  ParsedClause.setVarListDetails(VisitVarList(C.getVarList()),
+                                 /*IsReadOnly=*/false, /*IsZero=*/false);
+  if (SemaRef.OpenACC().CheckDeclareClause(ParsedClause))
+    return;
+  NewClause = OpenACCCopyClause::Create(
+      SemaRef.getASTContext(), ParsedClause.getClauseKind(),
+      ParsedClause.getBeginLoc(), ParsedClause.getLParenLoc(),
+      ParsedClause.getVarList(), ParsedClause.getEndLoc());
+}
+
+void OpenACCDeclClauseInstantiator::VisitLinkClause(
+    const OpenACCLinkClause &C) {
+  ParsedClause.setVarListDetails(
+      SemaRef.OpenACC().CheckLinkClauseVarList(VisitVarList(C.getVarList())),
+      /*IsReadOnly=*/false, /*IsZero=*/false);
+
+  if (SemaRef.OpenACC().CheckDeclareClause(ParsedClause))
+    return;
+
+  NewClause = OpenACCLinkClause::Create(
+      SemaRef.getASTContext(), ParsedClause.getBeginLoc(),
+      ParsedClause.getLParenLoc(), ParsedClause.getVarList(),
+      ParsedClause.getEndLoc());
+}
+
+void OpenACCDeclClauseInstantiator::VisitDeviceResidentClause(
+    const OpenACCDeviceResidentClause &C) {
+  ParsedClause.setVarListDetails(VisitVarList(C.getVarList()),
+                                 /*IsReadOnly=*/false, /*IsZero=*/false);
+  if (SemaRef.OpenACC().CheckDeclareClause(ParsedClause))
+    return;
+  NewClause = OpenACCDeviceResidentClause::Create(
+      SemaRef.getASTContext(), ParsedClause.getBeginLoc(),
+      ParsedClause.getLParenLoc(), ParsedClause.getVarList(),
+      ParsedClause.getEndLoc());
+}
+
+void OpenACCDeclClauseInstantiator::VisitCopyInClause(
+    const OpenACCCopyInClause &C) {
+  ParsedClause.setVarListDetails(VisitVarList(C.getVarList()), C.isReadOnly(),
+                                 /*IsZero=*/false);
+
+  if (SemaRef.OpenACC().CheckDeclareClause(ParsedClause))
+    return;
+  NewClause = OpenACCCopyInClause::Create(
+      SemaRef.getASTContext(), ParsedClause.getClauseKind(),
+      ParsedClause.getBeginLoc(), ParsedClause.getLParenLoc(),
+      ParsedClause.isReadOnly(), ParsedClause.getVarList(),
+      ParsedClause.getEndLoc());
+}
+void OpenACCDeclClauseInstantiator::VisitCopyOutClause(
+    const OpenACCCopyOutClause &C) {
+  ParsedClause.setVarListDetails(VisitVarList(C.getVarList()),
+                                 /*IsReadOnly=*/false, C.isZero());
+
+  if (SemaRef.OpenACC().CheckDeclareClause(ParsedClause))
+    return;
+  NewClause = OpenACCCopyOutClause::Create(
+      SemaRef.getASTContext(), ParsedClause.getClauseKind(),
+      ParsedClause.getBeginLoc(), ParsedClause.getLParenLoc(),
+      ParsedClause.isZero(), ParsedClause.getVarList(),
+      ParsedClause.getEndLoc());
+}
+void OpenACCDeclClauseInstantiator::VisitCreateClause(
+    const OpenACCCreateClause &C) {
+  ParsedClause.setVarListDetails(VisitVarList(C.getVarList()),
+                                 /*IsReadOnly=*/false, C.isZero());
+
+  if (SemaRef.OpenACC().CheckDeclareClause(ParsedClause))
+    return;
+  NewClause = OpenACCCreateClause::Create(
+      SemaRef.getASTContext(), ParsedClause.getClauseKind(),
+      ParsedClause.getBeginLoc(), ParsedClause.getLParenLoc(),
+      ParsedClause.isZero(), ParsedClause.getVarList(),
+      ParsedClause.getEndLoc());
+}
+void OpenACCDeclClauseInstantiator::VisitPresentClause(
+    const OpenACCPresentClause &C) {
+  ParsedClause.setVarListDetails(VisitVarList(C.getVarList()),
+                                 /*IsReadOnly=*/false, /*IsZero=*/false);
+  if (SemaRef.OpenACC().CheckDeclareClause(ParsedClause))
+    return;
+  NewClause = OpenACCPresentClause::Create(
+      SemaRef.getASTContext(), ParsedClause.getBeginLoc(),
+      ParsedClause.getLParenLoc(), ParsedClause.getVarList(),
+      ParsedClause.getEndLoc());
+}
+void OpenACCDeclClauseInstantiator::VisitDevicePtrClause(
+    const OpenACCDevicePtrClause &C) {
+  llvm::SmallVector<Expr *> VarList = VisitVarList(C.getVarList());
+  // Ensure each var is a pointer type.
+  VarList.erase(std::remove_if(VarList.begin(), VarList.end(),
+                               [&](Expr *E) {
+                                 return SemaRef.OpenACC().CheckVarIsPointerType(
+                                     OpenACCClauseKind::DevicePtr, E);
+                               }),
+                VarList.end());
+  ParsedClause.setVarListDetails(VarList,
+                                 /*IsReadOnly=*/false, /*IsZero=*/false);
+  if (SemaRef.OpenACC().CheckDeclareClause(ParsedClause))
+    return;
+  NewClause = OpenACCDevicePtrClause::Create(
+      SemaRef.getASTContext(), ParsedClause.getBeginLoc(),
+      ParsedClause.getLParenLoc(), ParsedClause.getVarList(),
+      ParsedClause.getEndLoc());
+}
+
+llvm::SmallVector<OpenACCClause *> InstantiateOpenACCClauseList(
+    Sema &S, const MultiLevelTemplateArgumentList &MLTAL,
+    OpenACCDirectiveKind DK, ArrayRef<const OpenACCClause *> ClauseList) {
+  llvm::SmallVector<OpenACCClause *> TransformedClauses;
+
+  for (const auto *Clause : ClauseList) {
+    SemaOpenACC::OpenACCParsedClause ParsedClause(DK, Clause->getClauseKind(),
+                                                  Clause->getBeginLoc());
+    ParsedClause.setEndLoc(Clause->getEndLoc());
+    if (const auto *WithParms = dyn_cast<OpenACCClauseWithParams>(Clause))
+      ParsedClause.setLParenLoc(WithParms->getLParenLoc());
+
+    OpenACCDeclClauseInstantiator Instantiator{S, MLTAL, ParsedClause};
+    Instantiator.Visit(Clause);
+    if (Instantiator.CreatedClause())
+      TransformedClauses.push_back(Instantiator.CreatedClause());
+  }
+  return TransformedClauses;
+}
+
+} // namespace
+
+Decl *TemplateDeclInstantiator::VisitOpenACCDeclareDecl(OpenACCDeclareDecl *D) {
+  SemaRef.OpenACC().ActOnConstruct(D->getDirectiveKind(), D->getBeginLoc());
+  llvm::SmallVector<OpenACCClause *> TransformedClauses =
+      InstantiateOpenACCClauseList(SemaRef, TemplateArgs, D->getDirectiveKind(),
+                                   D->clauses());
+
+  if (SemaRef.OpenACC().ActOnStartDeclDirective(D->getDirectiveKind(),
+                                                D->getBeginLoc()))
+    return nullptr;
+
+  DeclGroupRef Res = SemaRef.OpenACC().ActOnEndDeclDirective(
+      D->getDirectiveKind(), D->getBeginLoc(), D->getDirectiveLoc(),
+      D->getEndLoc(), TransformedClauses);
+
+  if (Res.isNull())
+    return nullptr;
+
+  return Res.getSingleDecl();
+}
+
 Decl *
 TemplateDeclInstantiator::VisitNamespaceAliasDecl(NamespaceAliasDecl *D) {
   NamespaceAliasDecl *Inst

diff  --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index eaabfae2409f4..ba56179cad7b1 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -11633,7 +11633,8 @@ class OpenACCClauseTransform final
       if (!Res.isUsable())
         continue;
 
-      Res = Self.getSema().OpenACC().ActOnVar(ParsedClause.getClauseKind(),
+      Res = Self.getSema().OpenACC().ActOnVar(ParsedClause.getDirectiveKind(),
+                                              ParsedClause.getClauseKind(),
                                               Res.get());
 
       if (Res.isUsable())
@@ -11698,7 +11699,8 @@ void OpenACCClauseTransform<Derived>::VisitSelfClause(
       if (!Res.isUsable())
         continue;
 
-      Res = Self.getSema().OpenACC().ActOnVar(ParsedClause.getClauseKind(),
+      Res = Self.getSema().OpenACC().ActOnVar(ParsedClause.getDirectiveKind(),
+                                              ParsedClause.getClauseKind(),
                                               Res.get());
 
       if (Res.isUsable())
@@ -11844,6 +11846,18 @@ void OpenACCClauseTransform<Derived>::VisitCopyClause(
       ParsedClause.getVarList(), ParsedClause.getEndLoc());
 }
 
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitLinkClause(
+    const OpenACCLinkClause &C) {
+  llvm_unreachable("link clause not valid unless a decl transform");
+}
+
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitDeviceResidentClause(
+    const OpenACCDeviceResidentClause &C) {
+  llvm_unreachable("device_resident clause not valid unless a decl transform");
+}
+
 template <typename Derived>
 void OpenACCClauseTransform<Derived>::VisitCopyInClause(
     const OpenACCCopyInClause &C) {

diff  --git a/clang/lib/Serialization/ASTCommon.cpp b/clang/lib/Serialization/ASTCommon.cpp
index 3a62c4ea5595b..ae8d0a61412b4 100644
--- a/clang/lib/Serialization/ASTCommon.cpp
+++ b/clang/lib/Serialization/ASTCommon.cpp
@@ -458,6 +458,7 @@ bool serialization::isRedeclarableDeclKind(unsigned Kind) {
   case Decl::RequiresExprBody:
   case Decl::UnresolvedUsingIfExists:
   case Decl::HLSLBuffer:
+  case Decl::OpenACCDeclare:
     return false;
 
   // These indirectly derive from Redeclarable<T> but are not actually

diff  --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
index 294e8e063e0a3..ea4b233a6c573 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -12515,7 +12515,7 @@ SmallVector<Expr *> ASTRecordReader::readOpenACCVarList() {
   unsigned NumVars = readInt();
   llvm::SmallVector<Expr *> VarList;
   for (unsigned I = 0; I < NumVars; ++I)
-    VarList.push_back(readSubExpr());
+    VarList.push_back(readExpr());
   return VarList;
 }
 
@@ -12781,10 +12781,20 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
     return OpenACCVectorClause::Create(getContext(), BeginLoc, LParenLoc,
                                        VectorExpr, EndLoc);
   }
+  case OpenACCClauseKind::Link: {
+    SourceLocation LParenLoc = readSourceLocation();
+    llvm::SmallVector<Expr *> VarList = readOpenACCVarList();
+    return OpenACCLinkClause::Create(getContext(), BeginLoc, LParenLoc, VarList,
+                                     EndLoc);
+  }
+  case OpenACCClauseKind::DeviceResident: {
+    SourceLocation LParenLoc = readSourceLocation();
+    llvm::SmallVector<Expr *> VarList = readOpenACCVarList();
+    return OpenACCDeviceResidentClause::Create(getContext(), BeginLoc,
+                                               LParenLoc, VarList, EndLoc);
+  }
 
   case OpenACCClauseKind::NoHost:
-  case OpenACCClauseKind::DeviceResident:
-  case OpenACCClauseKind::Link:
   case OpenACCClauseKind::Bind:
   case OpenACCClauseKind::Invalid:
     llvm_unreachable("Clause serialization not yet implemented");

diff  --git a/clang/lib/Serialization/ASTReaderDecl.cpp b/clang/lib/Serialization/ASTReaderDecl.cpp
index 17a41fff2267c..2a580c44b94e5 100644
--- a/clang/lib/Serialization/ASTReaderDecl.cpp
+++ b/clang/lib/Serialization/ASTReaderDecl.cpp
@@ -414,6 +414,8 @@ class ASTDeclReader : public DeclVisitor<ASTDeclReader, void> {
   void VisitEmptyDecl(EmptyDecl *D);
   void VisitLifetimeExtendedTemporaryDecl(LifetimeExtendedTemporaryDecl *D);
 
+  void VisitOpenACCDeclareDecl(OpenACCDeclareDecl *D);
+
   void VisitDeclContext(DeclContext *DC, uint64_t &LexicalOffset,
                         uint64_t &VisibleOffset, uint64_t &ModuleLocalOffset,
                         uint64_t &TULocalOffset);
@@ -3099,6 +3101,14 @@ void ASTDeclReader::VisitOMPCapturedExprDecl(OMPCapturedExprDecl *D) {
   VisitVarDecl(D);
 }
 
+void ASTDeclReader::VisitOpenACCDeclareDecl(OpenACCDeclareDecl *D) {
+  VisitDecl(D);
+  D->DirKind = Record.readEnum<OpenACCDirectiveKind>();
+  D->DirectiveLoc = Record.readSourceLocation();
+  D->EndLoc = Record.readSourceLocation();
+  Record.readOpenACCClauseList(D->Clauses);
+}
+
 //===----------------------------------------------------------------------===//
 // Attribute Reading
 //===----------------------------------------------------------------------===//
@@ -4204,6 +4214,9 @@ Decl *ASTReader::ReadDeclRecord(GlobalDeclID ID) {
     D = ImplicitConceptSpecializationDecl::CreateDeserialized(Context, ID,
                                                               Record.readInt());
     break;
+  case DECL_OPENACC_DECLARE:
+    D = OpenACCDeclareDecl::CreateDeserialized(Context, ID, Record.readInt());
+    break;
   }
 
   assert(D && "Unknown declaration reading AST file");

diff  --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index a01e72f193cf1..e51453881a94f 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -1135,6 +1135,7 @@ void ASTWriter::WriteBlockInfoBlock() {
   RECORD(DECL_OMP_DECLARE_REDUCTION);
   RECORD(DECL_OMP_ALLOCATE);
   RECORD(DECL_HLSL_BUFFER);
+  RECORD(DECL_OPENACC_DECLARE);
 
   // Statements and Exprs can occur in the Decls and Types block.
   AddStmtsExprs(Stream, Record);
@@ -8828,10 +8829,20 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
       AddStmt(const_cast<Expr *>(VC->getIntExpr()));
     return;
   }
+  case OpenACCClauseKind::Link: {
+    const auto *LC = cast<OpenACCLinkClause>(C);
+    writeSourceLocation(LC->getLParenLoc());
+    writeOpenACCVarList(LC);
+    return;
+  }
+  case OpenACCClauseKind::DeviceResident: {
+    const auto *DRC = cast<OpenACCDeviceResidentClause>(C);
+    writeSourceLocation(DRC->getLParenLoc());
+    writeOpenACCVarList(DRC);
+    return;
+  }
 
   case OpenACCClauseKind::NoHost:
-  case OpenACCClauseKind::DeviceResident:
-  case OpenACCClauseKind::Link:
   case OpenACCClauseKind::Bind:
   case OpenACCClauseKind::Invalid:
     llvm_unreachable("Clause serialization not yet implemented");

diff  --git a/clang/lib/Serialization/ASTWriterDecl.cpp b/clang/lib/Serialization/ASTWriterDecl.cpp
index ac80bb46afa2d..a1810003f5425 100644
--- a/clang/lib/Serialization/ASTWriterDecl.cpp
+++ b/clang/lib/Serialization/ASTWriterDecl.cpp
@@ -176,6 +176,8 @@ namespace clang {
     void VisitOMPDeclareMapperDecl(OMPDeclareMapperDecl *D);
     void VisitOMPCapturedExprDecl(OMPCapturedExprDecl *D);
 
+    void VisitOpenACCDeclareDecl(OpenACCDeclareDecl *D);
+
     /// Add an Objective-C type parameter list to the given record.
     void AddObjCTypeParamList(ObjCTypeParamList *typeParams) {
       // Empty type parameter list.
@@ -2258,6 +2260,16 @@ void ASTDeclWriter::VisitOMPCapturedExprDecl(OMPCapturedExprDecl *D) {
   Code = serialization::DECL_OMP_CAPTUREDEXPR;
 }
 
+void ASTDeclWriter::VisitOpenACCDeclareDecl(OpenACCDeclareDecl *D) {
+  Record.writeUInt32(D->clauses().size());
+  VisitDecl(D);
+  Record.writeEnum(D->DirKind);
+  Record.AddSourceLocation(D->DirectiveLoc);
+  Record.AddSourceLocation(D->EndLoc);
+  Record.writeOpenACCClauseList(D->clauses());
+  Code = serialization::DECL_OPENACC_DECLARE;
+}
+
 //===----------------------------------------------------------------------===//
 // ASTWriter Implementation
 //===----------------------------------------------------------------------===//

diff  --git a/clang/test/AST/ast-print-openacc-declare-construct.cpp b/clang/test/AST/ast-print-openacc-declare-construct.cpp
new file mode 100644
index 0000000000000..fce4afc6aedae
--- /dev/null
+++ b/clang/test/AST/ast-print-openacc-declare-construct.cpp
@@ -0,0 +1,52 @@
+// RUN: %clang_cc1 -fopenacc -ast-print %s -o - | FileCheck %s
+
+int *Global, *Global2;
+int GlobalArray[5];
+int GlobalArray2[5];
+// CHECK: #pragma acc declare deviceptr(Global) copyin(GlobalArray)
+#pragma acc declare deviceptr(Global), copyin(GlobalArray)
+// CHECK: #pragma acc declare create(Global2, GlobalArray2)
+#pragma acc declare create(Global2, GlobalArray2)
+
+namespace NS {
+int NSVar;
+int NSArray[5];
+// CHECK: #pragma acc declare create(NSVar, NSArray)
+#pragma acc declare create(NSVar, NSArray)
+}
+
+struct Struct {
+  static const int StaticMem = 5;
+  static const int StaticMemArray[5];
+// CHECK: #pragma acc declare copyin(StaticMem, StaticMemArray)
+#pragma acc declare copyin(StaticMem, StaticMemArray)
+
+  void MemFunc1(int Arg) {
+    int Local;
+    int LocalArray[5];
+// CHECK: #pragma acc declare present(Arg, Local, LocalArray)
+#pragma acc declare present(Arg, Local, LocalArray)
+  }
+  void MemFunc2(int Arg);
+};
+void Struct::MemFunc2(int Arg) {
+  int Local;
+  int LocalArray[5];
+// CHECK: #pragma acc declare present(Arg, Local, LocalArray)
+#pragma acc declare present(Arg, Local, LocalArray)
+}
+
+void NormalFunc(int Arg) {
+  int Local;
+  int LocalArray[5];
+// CHECK: #pragma acc declare present(Arg, Local, LocalArray)
+#pragma acc declare present(Arg, Local, LocalArray)
+}
+
+void NormalFunc2(int *Arg) {
+  int Local;
+  int LocalArray[5];
+  extern int ExternLocal;
+// CHECK: #pragma acc declare deviceptr(Arg) device_resident(Local) link(ExternLocal)
+#pragma acc declare deviceptr(Arg) device_resident(Local) link(ExternLocal)
+}

diff  --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c
index b871624b6e943..7c987ba0a1bd7 100644
--- a/clang/test/ParserOpenACC/parse-clauses.c
+++ b/clang/test/ParserOpenACC/parse-clauses.c
@@ -520,23 +520,25 @@ 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 'device_resident' not yet implemented, clause ignored}}
-#pragma acc serial device_resident(s.array[s.value] s.array[s.value :5] ), self
-  for(int i = 0; i < 5;++i) {}
+  // expected-error at +3{{expected ','}}
+  // expected-error at +2{{OpenACC variable on 'declare' construct is not a valid variable name or array name}}
+  // expected-error at +1{{OpenACC variable on 'declare' construct is not a valid variable name or array name}}
+#pragma acc declare device_resident(s.array[s.value] s.array[s.value :5] ), copy(s)
 
-  // expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented, clause ignored}}
-#pragma acc serial device_resident(s.array[s.value : 5], s.value), self
-  for(int i = 0; i < 5;++i) {}
+  int CopyRef1, CopyRef2, CopyRef3;
 
-  // expected-error at +2{{expected ','}}
-  // expected-warning at +1{{OpenACC clause 'link' not yet implemented, clause ignored}}
-#pragma acc serial link(s.array[s.value] s.array[s.value :5] ), self
-  for(int i = 0; i < 5;++i) {}
+  // expected-error at +2{{OpenACC variable on 'declare' construct is not a valid variable name or array name}}
+  // expected-error at +1{{OpenACC variable on 'declare' construct is not a valid variable name or array name}}
+#pragma acc declare device_resident(s.array[s.value : 5], s.value), copy(CopyRef1)
 
-  // expected-warning at +1{{OpenACC clause 'link' not yet implemented, clause ignored}}
-#pragma acc serial link(s.array[s.value : 5], s.value), self
-  for(int i = 0; i < 5;++i) {}
+  // expected-error at +3{{expected ','}}
+  // expected-error at +2{{OpenACC variable on 'declare' construct is not a valid variable name or array name}}
+  // expected-error at +1{{OpenACC variable on 'declare' construct is not a valid variable name or array name}}
+#pragma acc declare link(s.array[s.value] s.array[s.value :5] ), copy(CopyRef2)
+
+  // expected-error at +2{{OpenACC variable on 'declare' construct is not a valid variable name or array name}}
+  // expected-error at +1{{OpenACC variable on 'declare' construct is not a valid variable name or array name}}
+#pragma acc declare link(s.array[s.value : 5], s.value), copy(CopyRef3)
 
   // expected-error at +1{{expected ','}}
 #pragma acc update host(s.array[s.value] s.array[s.value :5] )

diff  --git a/clang/test/ParserOpenACC/parse-constructs.c b/clang/test/ParserOpenACC/parse-constructs.c
index f0698495a3cc2..976220979f0b9 100644
--- a/clang/test/ParserOpenACC/parse-constructs.c
+++ b/clang/test/ParserOpenACC/parse-constructs.c
@@ -131,7 +131,7 @@ void func() {
 
 
   // expected-error at +2{{invalid OpenACC clause 'clause'}}
-  // expected-warning at +1{{OpenACC construct 'declare' not yet implemented, pragma ignored}}
+  // expected-error at +1{{no valid clauses specified in OpenACC 'declare' directive}}
 #pragma acc declare clause list
   for(;;){}
   // expected-error at +1{{invalid OpenACC clause 'clause'}}

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 9e74ce27ffbd9..bbc40f4b401b2 100644
--- a/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
+++ b/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
@@ -80,7 +80,7 @@ void uses() {
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop auto deviceptr(VarPtr)
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented}}
+  // expected-error at +1{{OpenACC 'device_resident' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop auto device_resident(VarPtr)
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop auto firstprivate(Var)
@@ -88,7 +88,7 @@ void uses() {
   // expected-error at +1{{OpenACC 'host' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop auto host(Var)
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning at +1{{OpenACC clause 'link' not yet implemented}}
+  // expected-error at +1{{OpenACC 'link' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop auto link(Var)
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop auto no_create(Var)
@@ -197,7 +197,7 @@ void uses() {
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop deviceptr(VarPtr) auto
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented}}
+  // expected-error at +1{{OpenACC 'device_resident' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop device_resident(VarPtr) auto
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop firstprivate(Var) auto
@@ -205,7 +205,7 @@ void uses() {
   // expected-error at +1{{OpenACC 'host' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop host(Var) auto
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning at +1{{OpenACC clause 'link' not yet implemented}}
+  // expected-error at +1{{OpenACC 'link' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop link(Var) auto
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop no_create(Var) auto
@@ -315,7 +315,7 @@ void uses() {
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop independent deviceptr(VarPtr)
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented}}
+  // expected-error at +1{{OpenACC 'device_resident' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop independent device_resident(VarPtr)
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop independent firstprivate(Var)
@@ -323,7 +323,7 @@ void uses() {
   // expected-error at +1{{OpenACC 'host' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop independent host(Var)
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning at +1{{OpenACC clause 'link' not yet implemented}}
+  // expected-error at +1{{OpenACC 'link' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop independent link(Var)
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop independent no_create(Var)
@@ -432,7 +432,7 @@ void uses() {
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop deviceptr(VarPtr) independent
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented}}
+  // expected-error at +1{{OpenACC 'device_resident' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop device_resident(VarPtr) independent
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop firstprivate(Var) independent
@@ -440,7 +440,7 @@ void uses() {
   // expected-error at +1{{OpenACC 'host' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop host(Var) independent
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning at +1{{OpenACC clause 'link' not yet implemented}}
+  // expected-error at +1{{OpenACC 'link' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop link(Var) independent
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop no_create(Var) independent
@@ -558,7 +558,7 @@ void uses() {
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop seq deviceptr(VarPtr)
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented}}
+  // expected-error at +1{{OpenACC 'device_resident' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop seq device_resident(VarPtr)
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop seq firstprivate(Var)
@@ -566,7 +566,7 @@ void uses() {
   // expected-error at +1{{OpenACC 'host' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop seq host(Var)
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning at +1{{OpenACC clause 'link' not yet implemented}}
+  // expected-error at +1{{OpenACC 'link' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop seq link(Var)
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop seq no_create(Var)
@@ -681,7 +681,7 @@ void uses() {
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop deviceptr(VarPtr) seq
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented}}
+  // expected-error at +1{{OpenACC 'device_resident' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop device_resident(VarPtr) seq
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop firstprivate(Var) seq
@@ -689,7 +689,7 @@ void uses() {
   // expected-error at +1{{OpenACC 'host' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop host(Var) seq
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning at +1{{OpenACC clause 'link' not yet implemented}}
+  // expected-error at +1{{OpenACC 'link' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop link(Var) seq
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop no_create(Var) seq

diff  --git a/clang/test/SemaOpenACC/combined-construct-device_type-clause.c b/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
index c185f607548ff..2d7a9418cce15 100644
--- a/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
@@ -107,8 +107,7 @@ void uses() {
   // expected-note at +1{{previous clause is here}}
 #pragma acc serial loop device_type(*) deviceptr(VarPtr)
   for(int i = 0; i < 5; ++i);
-  // expected-error at +2{{OpenACC clause 'device_resident' 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 'device_resident' clause is not valid on 'kernels loop' directive}}
 #pragma acc kernels loop device_type(*)  device_resident(VarPtr)
   for(int i = 0; i < 5; ++i);
   // expected-error at +2{{OpenACC clause 'firstprivate' may not follow a 'device_type' clause in a 'parallel loop' construct}}
@@ -118,8 +117,7 @@ void uses() {
   // expected-error at +1{{OpenACC 'host' clause is not valid on 'serial loop' directive}}
 #pragma acc serial loop device_type(*) host(Var)
   for(int i = 0; i < 5; ++i);
-  // expected-error at +2{{OpenACC clause 'link' may not follow a 'device_type' clause in a 'parallel loop' construct}}
-  // expected-note at +1{{previous clause is here}}
+  // expected-error at +1{{OpenACC 'link' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop device_type(*) link(Var)
   for(int i = 0; i < 5; ++i);
   // expected-error at +2{{OpenACC clause 'no_create' may not follow a 'device_type' clause in a 'serial 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 4290fb7665685..5f464cc94b092 100644
--- a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
@@ -113,8 +113,7 @@ void uses() {
   // expected-note at +1{{previous clause is here}}
 #pragma acc kernels device_type(*) deviceptr(VarPtr)
   while(1);
-  // expected-error at +2{{OpenACC clause 'device_resident' may not follow a 'device_type' clause in a 'kernels' construct}}
-  // expected-note at +1{{previous clause is here}}
+  // expected-error at +1{{OpenACC 'device_resident' clause is not valid on 'kernels' directive}}
 #pragma acc kernels device_type(*)  device_resident(VarPtr)
   while(1);
   // expected-error at +2{{OpenACC clause 'firstprivate' may not follow a 'device_type' clause in a 'parallel' construct}}
@@ -124,8 +123,7 @@ void uses() {
   // expected-error at +1{{OpenACC 'host' clause is not valid on 'kernels' directive}}
 #pragma acc kernels device_type(*) host(Var)
   while(1);
-  // expected-error at +2{{OpenACC clause 'link' may not follow a 'device_type' clause in a 'kernels' construct}}
-  // expected-note at +1{{previous clause is here}}
+  // expected-error at +1{{OpenACC 'link' clause is not valid on 'kernels' directive}}
 #pragma acc kernels device_type(*) link(Var)
   while(1);
   // expected-error at +2{{OpenACC clause 'no_create' may not follow a 'device_type' clause in a 'kernels' construct}}

diff  --git a/clang/test/SemaOpenACC/data-construct.cpp b/clang/test/SemaOpenACC/data-construct.cpp
index 4c868b68e332e..e1d0c04d7fdee 100644
--- a/clang/test/SemaOpenACC/data-construct.cpp
+++ b/clang/test/SemaOpenACC/data-construct.cpp
@@ -165,7 +165,6 @@ struct HasMembers {
   // 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-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)
   ;
   }

diff  --git a/clang/test/SemaOpenACC/declare-construct-ast.cpp b/clang/test/SemaOpenACC/declare-construct-ast.cpp
new file mode 100644
index 0000000000000..dfc79039fc9e6
--- /dev/null
+++ b/clang/test/SemaOpenACC/declare-construct-ast.cpp
@@ -0,0 +1,390 @@
+// 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;
+// CHECK: VarDecl{{.*}}Global 'int *'
+int GlobalArray[5];
+// CHECK-NEXT: VarDecl{{.*}}GlobalArray 'int[5]'
+#pragma acc declare deviceptr(Global), copyin(GlobalArray)
+// CHECK-NEXT: OpenACCDeclareDecl
+// CHECK-NEXT: deviceptr clause
+// CHECK-NEXT: DeclRefExpr{{.*}}'Global' 'int *'
+// CHECK-NEXT: copyin clause
+// CHECK-NEXT: DeclRefExpr{{.*}}'GlobalArray' 'int[5]'
+
+int *Global2;
+// CHECK: VarDecl{{.*}}Global2 'int *'
+int GlobalArray2[5];
+// CHECK-NEXT: VarDecl{{.*}}GlobalArray2 'int[5]'
+#pragma acc declare create(Global2, GlobalArray2)
+// CHECK-NEXT: OpenACCDeclareDecl
+// CHECK-NEXT: create clause
+// CHECK-NEXT: DeclRefExpr{{.*}}'Global2' 'int *'
+// CHECK-NEXT: DeclRefExpr{{.*}}'GlobalArray2' 'int[5]'
+
+int Global3;
+// CHECK: VarDecl{{.*}}Global3 'int'
+int GlobalArray3[5];
+// CHECK-NEXT: VarDecl{{.*}}GlobalArray3 'int[5]'
+#pragma acc declare link(Global3) device_resident(GlobalArray3)
+// CHECK-NEXT: OpenACCDeclareDecl
+// CHECK-NEXT: link clause
+// CHECK-NEXT: DeclRefExpr{{.*}}'Global3' 'int'
+// CHECK-NEXT: device_resident clause
+// CHECK-NEXT: DeclRefExpr{{.*}}'GlobalArray3' 'int[5]'
+
+namespace NS {
+int NSVar;
+// CHECK: VarDecl{{.*}}NSVar 'int'
+int NSArray[5];
+// CHECK-NEXT: VarDecl{{.*}}NSArray 'int[5]'
+#pragma acc declare create(NSVar, NSArray)
+// CHECK-NEXT: OpenACCDeclareDecl
+// CHECK-NEXT: create clause
+// CHECK-NEXT: DeclRefExpr{{.*}}'NSVar' 'int'
+// CHECK-NEXT: DeclRefExpr{{.*}}'NSArray' 'int[5]'
+}
+
+struct Struct {
+  // CHECK-NEXT: CXXRecordDecl{{.*}} Struct definition
+  // Skip DefinitionData and go right to the definition.
+  // CHECK: CXXRecordDecl{{.*}} implicit struct Struct
+  static const int StaticMem = 5;
+  // CHECK-NEXT: VarDecl {{.*}} StaticMem 'const int' static cinit
+  // CHECK-NEXT: IntegerLiteral{{.*}}'int' 5
+  static const int StaticMem2 = 5;
+  // CHECK-NEXT: VarDecl {{.*}} StaticMem2 'const int' static cinit
+  // CHECK-NEXT: IntegerLiteral{{.*}}'int' 5
+  static const int StaticMemArray[5];
+  // CHECK-NEXT: VarDecl {{.*}} StaticMemArray 'const int[5]' static
+  static const int StaticMemArray2[5];
+  // CHECK-NEXT: VarDecl {{.*}} StaticMemArray2 'const int[5]' static
+#pragma acc declare copyin(StaticMem, StaticMemArray) create(StaticMem2, StaticMemArray2)
+  // CHECK-NEXT: OpenACCDeclareDecl
+  // CHECK-NEXT: copyin clause
+  // CHECK-NEXT: DeclRefExpr{{.*}}'StaticMem' 'const int'
+  // CHECK-NEXT: DeclRefExpr{{.*}}'StaticMemArray' 'const int[5]'
+  // CHECK-NEXT: create clause
+  // CHECK-NEXT: DeclRefExpr{{.*}}'StaticMem2' 'const int'
+  // CHECK-NEXT: DeclRefExpr{{.*}}'StaticMemArray2' 'const int[5]'
+
+  void MemFunc1(int Arg) {
+    // CHECK-NEXT: CXXMethodDecl{{.*}}MemFunc1 'void (int)'
+    // CHECK-NEXT: ParmVarDecl{{.*}} Arg 'int'
+    // CHECK-NEXT: CompoundStmt
+    int Local;
+    // CHECK-NEXT: DeclStmt
+    // CHECK-NEXT: VarDecl{{.*}} Local 'int'
+    int LocalArray[5];
+    // CHECK-NEXT: DeclStmt
+    // CHECK-NEXT: VarDecl{{.*}} LocalArray 'int[5]'
+#pragma acc declare present(Arg, Local, LocalArray)
+    // CHECK-NEXT: DeclStmt
+    // CHECK-NEXT: OpenACCDeclareDecl
+    // CHECK-NEXT: present clause
+    // CHECK-NEXT: DeclRefExpr{{.*}}'Arg' 'int'
+    // CHECK-NEXT: DeclRefExpr{{.*}}'Local' 'int'
+    // CHECK-NEXT: DeclRefExpr{{.*}}'LocalArray' 'int[5]'
+  }
+  void MemFunc2(int Arg);
+  // CHECK: CXXMethodDecl{{.*}}MemFunc2
+};
+void Struct::MemFunc2(int Arg) {
+  // CHECK: CXXMethodDecl{{.*}}MemFunc2 'void (int)'
+  // CHECK-NEXT: ParmVarDecl{{.*}} Arg 'int'
+  // CHECK-NEXT: CompoundStmt
+  int Local;
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} Local 'int'
+  int LocalArray[5];
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} LocalArray 'int[5]'
+#pragma acc declare present(Arg, Local, LocalArray)
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: OpenACCDeclareDecl
+  // CHECK-NEXT: present clause
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Arg' 'int'
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Local' 'int'
+  // CHECK-NEXT: DeclRefExpr{{.*}}'LocalArray' 'int[5]'
+}
+
+void NormalFunc(int Arg) {
+  // CHECK-NEXT: FunctionDecl{{.*}}NormalFunc 'void (int)'
+  // CHECK-NEXT: ParmVarDecl{{.*}} Arg 'int'
+  // CHECK-NEXT: CompoundStmt
+  int Local;
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} Local 'int'
+  int LocalArray[5];
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} LocalArray 'int[5]'
+#pragma acc declare present(Arg, Local, LocalArray)
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: OpenACCDeclareDecl
+  // CHECK-NEXT: present clause
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Arg' 'int'
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Local' 'int'
+  // CHECK-NEXT: DeclRefExpr{{.*}}'LocalArray' 'int[5]'
+}
+
+template<typename T>
+struct DependentStruct {
+  // CHECK: ClassTemplateDecl{{.*}}DependentStruct
+  // CHECK-NEXT: TemplateTypeParmDecl{{.*}}depth 0 index 0 T
+  // CHECK-NEXT: CXXRecordDecl{{.*}}DependentStruct definition
+  // CHECK: CXXRecordDecl{{.*}}implicit struct DependentStruct
+  static const T StaticMem = 5;
+  // CHECK-NEXT: VarDecl{{.*}} StaticMem 'const T' static cinit
+  // CHECK-NEXT: IntegerLiteral{{.*}}'int' 5
+  static const T StaticMem2 = 5;
+  // CHECK-NEXT: VarDecl{{.*}} StaticMem2 'const T' static cinit
+  // CHECK-NEXT: IntegerLiteral{{.*}}'int' 5
+  static constexpr T StaticMemArray[5] = {};
+  // CHECK-NEXT: VarDecl{{.*}} StaticMemArray 'const T[5]'
+  // CHECK-NEXT: InitListExpr{{.*}}'void'
+  static constexpr T StaticMemArray2[5] = {};
+  // CHECK-NEXT: VarDecl{{.*}} StaticMemArray2 'const T[5]'
+  // CHECK-NEXT: InitListExpr{{.*}}'void'
+#pragma acc declare copyin(StaticMem, StaticMemArray) create(StaticMem2, StaticMemArray2)
+  // CHECK-NEXT: OpenACCDeclareDecl
+  // CHECK-NEXT: copyin clause
+  // CHECK-NEXT: DeclRefExpr{{.*}}'StaticMem' 'const T'
+  // CHECK-NEXT: DeclRefExpr{{.*}}'StaticMemArray' 'const T[5]'
+  // CHECK-NEXT: create clause
+  // CHECK-NEXT: DeclRefExpr{{.*}}'StaticMem2' 'const T'
+  // CHECK-NEXT: DeclRefExpr{{.*}}'StaticMemArray2' 'const T[5]'
+
+  template<typename U>
+  void DepMemFunc1(U Arg, U Arg2) {
+    // CHECK-NEXT: FunctionTemplateDecl{{.*}}DepMemFunc1
+    // CHECK-NEXT: TemplateTypeParmDecl{{.*}}depth 1 index 0 U
+    // CHECK-NEXT: CXXMethodDecl{{.*}}DepMemFunc1 'void (U, U)'
+    // CHECK-NEXT: ParmVarDecl{{.*}} Arg 'U'
+    // CHECK-NEXT: ParmVarDecl{{.*}} Arg2 'U'
+    // CHECK-NEXT: CompoundStmt
+    T Local, Local2;
+    // CHECK-NEXT: DeclStmt
+    // CHECK-NEXT: VarDecl{{.*}} Local 'T'
+    // CHECK-NEXT: VarDecl{{.*}} Local2 'T'
+    U LocalArray[5];
+    // CHECK-NEXT: DeclStmt
+    // CHECK-NEXT: VarDecl{{.*}} LocalArray 'U[5]'
+    U LocalArray2[5];
+    // CHECK-NEXT: DeclStmt
+    // CHECK-NEXT: VarDecl{{.*}} LocalArray2 'U[5]'
+#pragma acc declare copy(Arg, Local, LocalArray) copyout(Arg2, Local2, LocalArray2)
+    // CHECK-NEXT: DeclStmt
+    // CHECK-NEXT: OpenACCDeclareDecl
+    // CHECK-NEXT: copy clause
+    // CHECK-NEXT: DeclRefExpr{{.*}}'Arg' 'U'
+    // CHECK-NEXT: DeclRefExpr{{.*}}'Local' 'T'
+    // CHECK-NEXT: DeclRefExpr{{.*}}'LocalArray' 'U[5]'
+    // CHECK-NEXT: copyout clause
+    // CHECK-NEXT: DeclRefExpr{{.*}}'Arg2' 'U'
+    // CHECK-NEXT: DeclRefExpr{{.*}}'Local2' 'T'
+    // CHECK-NEXT: DeclRefExpr{{.*}}'LocalArray2' 'U[5]'
+
+    extern T Local3;
+    // CHECK-NEXT: DeclStmt
+    // CHECK-NEXT: VarDecl{{.*}} Local3 'T' extern
+    T Local4;
+    // CHECK-NEXT: DeclStmt
+    // CHECK-NEXT: VarDecl{{.*}} Local4 'T'
+#pragma acc declare link(Local3) device_resident(Local4)
+    // CHECK-NEXT: DeclStmt
+    // CHECK-NEXT: OpenACCDeclareDecl
+    // CHECK-NEXT: link clause
+    // CHECK-NEXT: DeclRefExpr{{.*}}'Local3' 'T'
+    // CHECK-NEXT: device_resident clause
+    // CHECK-NEXT: DeclRefExpr{{.*}}'Local4' 'T'
+  }
+  template<typename U>
+  void DepMemFunc2(U Arg);
+  // CHECK-NEXT: FunctionTemplateDecl{{.*}}DepMemFunc2
+  // CHECK-NEXT: TemplateTypeParmDecl{{.*}}depth 1 index 0 U
+  // CHECK-NEXT: CXXMethodDecl{{.*}}DepMemFunc2 'void (U)'
+  // CHECK-NEXT: ParmVarDecl{{.*}} Arg 'U'
+};
+
+// Instantiation of class.
+// CHECK-NEXT: ClassTemplateSpecializationDecl{{.*}}DependentStruct definition
+// CHECK: TemplateArgument type 'int'
+// CHECK-NEXT: BuiltinType{{.*}}'int'
+// CHECK-NEXT: CXXRecordDecl{{.*}} struct DependentStruct
+
+// CHECK-NEXT: VarDecl{{.*}} StaticMem 'const int' 
+// CHECK-NEXT: IntegerLiteral{{.*}}'int' 5
+
+// CHECK-NEXT: VarDecl{{.*}} StaticMem2 'const int' 
+// CHECK-NEXT: IntegerLiteral{{.*}}'int' 5
+//
+// CHECK-NEXT: VarDecl{{.*}} StaticMemArray 'const int[5]'
+// CHECK-NEXT: value: Array size=5
+// CHECK-NEXT: filler: 5 x Int 0
+// CHECK-NEXT: InitListExpr{{.*}} 'const int[5]'
+// CHECK-NEXT: array_filler
+
+// CHECK-NEXT: VarDecl{{.*}} StaticMemArray2 'const int[5]'
+// CHECK-NEXT: value: Array size=5
+// CHECK-NEXT: filler: 5 x Int 0
+// CHECK-NEXT: InitListExpr{{.*}} 'const int[5]'
+// CHECK-NEXT: array_filler
+
+// CHECK-NEXT: OpenACCDeclareDecl
+// CHECK-NEXT: copyin clause
+// CHECK-NEXT: DeclRefExpr{{.*}}'StaticMem' 'const int'
+// CHECK-NEXT: DeclRefExpr{{.*}}'StaticMemArray' 'const int[5]'
+// CHECK-NEXT: create clause
+// CHECK-NEXT: DeclRefExpr{{.*}}'StaticMem2' 'const int'
+// CHECK-NEXT: DeclRefExpr{{.*}}'StaticMemArray2' 'const int[5]'
+
+// CHECK-NEXT: FunctionTemplateDecl{{.*}} DepMemFunc1
+// CHECK-NEXT: TemplateTypeParmDecl{{.*}}depth 0 index 0 U
+// CHECK-NEXT: CXXMethodDecl{{.*}}DepMemFunc1 'void (U, U)'
+// CHECK-NEXT: ParmVarDecl{{.*}} Arg 'U'
+// CHECK-NEXT: ParmVarDecl{{.*}} Arg2 'U'
+// CHECK-NEXT: CXXMethodDecl{{.*}}DepMemFunc1 'void (float, float)'
+// CHECK-NEXT: TemplateArgument type 'float'
+// CHECK-NEXT: BuiltinType{{.*}}'float'
+// CHECK-NEXT: ParmVarDecl{{.*}} Arg 'float'
+// CHECK-NEXT: ParmVarDecl{{.*}} Arg2 'float'
+// CHECK-NEXT: CompoundStmt
+
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT: VarDecl{{.*}} Local 'int'
+// CHECK-NEXT: VarDecl{{.*}} Local2 'int'
+
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT: VarDecl{{.*}} LocalArray 'float[5]'
+
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT: VarDecl{{.*}} LocalArray2 'float[5]'
+
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT: OpenACCDeclareDecl
+// CHECK-NEXT: copy clause
+// CHECK-NEXT: DeclRefExpr{{.*}}'Arg' 'float'
+// CHECK-NEXT: DeclRefExpr{{.*}}'Local' 'int'
+// CHECK-NEXT: DeclRefExpr{{.*}}'LocalArray' 'float[5]'
+// CHECK-NEXT: copyout clause
+// CHECK-NEXT: DeclRefExpr{{.*}}'Arg2' 'float'
+// CHECK-NEXT: DeclRefExpr{{.*}}'Local2' 'int'
+// CHECK-NEXT: DeclRefExpr{{.*}}'LocalArray2' 'float[5]'
+
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT: VarDecl{{.*}} Local3 'int' extern
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT: VarDecl{{.*}} Local4 'int'
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT: OpenACCDeclareDecl
+// CHECK-NEXT: link clause
+// CHECK-NEXT: DeclRefExpr{{.*}}'Local3' 'int'
+// CHECK-NEXT: device_resident clause
+// CHECK-NEXT: DeclRefExpr{{.*}}'Local4' 'int'
+
+// CHECK-NEXT: FunctionTemplateDecl{{.*}}DepMemFunc2
+// CHECK-NEXT: TemplateTypeParmDecl{{.*}}depth 0 index 0 U
+// CHECK-NEXT: CXXMethodDecl{{.*}}DepMemFunc2 'void (U)'
+// CHECK-NEXT: ParmVarDecl{{.*}} Arg 'U'
+// CHECK-NEXT: CXXMethodDecl{{.*}}DepMemFunc2 'void (float)'
+// CHECK-NEXT: TemplateArgument type 'float'
+// CHECK-NEXT: BuiltinType{{.*}}'float'
+// CHECK-NEXT: ParmVarDecl{{.*}} Arg 'float'
+// CHECK-NEXT: CompoundStmt
+
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT: VarDecl{{.*}} Local 'int'
+
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT: VarDecl{{.*}} LocalArray 'float[5]'
+
+// CHECK-NEXT: DeclStmt
+// CHECK-NEXT: OpenACCDeclareDecl
+// CHECK-NEXT: present clause
+// CHECK-NEXT: DeclRefExpr{{.*}}'Arg' 'float'
+// CHECK-NEXT: DeclRefExpr{{.*}}'Local' 'int'
+// CHECK-NEXT: DeclRefExpr{{.*}}'LocalArray' 'float[5]'
+
+template<typename T>
+template<typename U>
+void DependentStruct<T>::DepMemFunc2(U Arg) {
+  // CHECK: FunctionTemplateDecl{{.*}} DepMemFunc2
+  // CHECK-NEXT: TemplateTypeParmDecl{{.*}}depth 1 index 0 U
+  // CHECK-NEXT: CXXMethodDecl{{.*}}DepMemFunc2 'void (U)'
+  // CHECK-NEXT: ParmVarDecl{{.*}} Arg 'U'
+  // CHECK-NEXT: CompoundStmt
+  T Local;
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} Local 'T'
+  U LocalArray[5];
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} LocalArray 'U[5]'
+#pragma acc declare present(Arg, Local, LocalArray)
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: OpenACCDeclareDecl
+  // CHECK-NEXT: present clause
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Arg' 'U'
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Local' 'T'
+  // CHECK-NEXT: DeclRefExpr{{.*}}'LocalArray' 'U[5]'
+}
+
+template<typename T, unsigned Size>
+void DependentFunc(T Arg) {
+  // CHECK: FunctionTemplateDecl{{.*}} DependentFunc
+  // CHECK-NEXT: TemplateTypeParmDecl{{.*}}depth 0 index 0 T
+  // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} depth 0 index 1 Size
+  // CHECK-NEXT: FunctionDecl{{.*}}DependentFunc 'void (T)'
+  // CHECK-NEXT: ParmVarDecl{{.*}} Arg 'T'
+  // CHECK-NEXT: CompoundStmt
+  T Local;
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} Local 'T'
+  T LocalArray[Size];
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} LocalArray 'T[Size]'
+
+#pragma acc declare present(Arg, Local, LocalArray)
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: OpenACCDeclareDecl
+  // CHECK-NEXT: present clause
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Arg' 'T'
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Local' 'T'
+  // CHECK-NEXT: DeclRefExpr{{.*}}'LocalArray' 'T[Size]'
+
+  // Instantiation:
+  // CHECK-NEXT: FunctionDecl{{.*}} DependentFunc 'void (int)'
+  // CHECK-NEXT: TemplateArgument type 'int'
+  // CHECK-NEXT: BuiltinType{{.*}}'int'
+  // CHECK-NEXT: TemplateArgument integral '5U'
+  // CHECK-NEXT: ParmVarDecl{{.*}} Arg 'int'
+  // CHECK-NEXT: CompoundStmt
+
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} Local 'int'
+
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} LocalArray 'int[5]'
+
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: OpenACCDeclareDecl
+  // CHECK-NEXT: present clause
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Arg' 'int'
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Local' 'int'
+  // CHECK-NEXT: DeclRefExpr{{.*}}'LocalArray' 'int[5]'
+}
+
+void use() {
+  float i;
+  DependentStruct<int> S;
+  S.DepMemFunc1(i, i);
+  S.DepMemFunc2(i);
+  DependentFunc<int, 5>(i);
+}
+
+#endif // PCH_HELPER

diff  --git a/clang/test/SemaOpenACC/declare-construct.cpp b/clang/test/SemaOpenACC/declare-construct.cpp
new file mode 100644
index 0000000000000..a1fed096635fa
--- /dev/null
+++ b/clang/test/SemaOpenACC/declare-construct.cpp
@@ -0,0 +1,308 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+int *Global;
+int GlobalArray[5];
+// expected-error at +1{{no valid clauses specified in OpenACC 'declare' directive}}
+#pragma acc declare
+namespace NS {
+int *NSVar;
+int NSArray[5];
+// expected-error at +2{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+// expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(Global, GlobalArray)
+// Ok, correct scope.
+#pragma acc declare create(NSVar, NSArray)
+
+// expected-error at +4{{variable referenced in 'create' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -3{{previous reference is here}}
+// expected-error at +2{{variable referenced in 'copyin' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -5{{previous reference is here}}
+#pragma acc declare create(NSVar) copyin(NSVar)
+
+// expected-error at +1{{no valid clauses specified in OpenACC 'declare' directive}}
+#pragma acc declare
+
+int NSVar1, NSVar2, NSVar3, NSVar4, NSVar5, *NSVar6, NSVar7, NSVar8;
+
+// Only create, copyin, deviceptr, device-resident, link at NS scope.
+// expected-error at +3{{OpenACC 'copy' clause on a 'declare' directive is not allowed at global or namespace scope}}
+// expected-error at +2{{OpenACC 'copyout' clause on a 'declare' directive is not allowed at global or namespace scope}}
+// expected-error at +1{{OpenACC 'present' clause on a 'declare' directive is not allowed at global or namespace scope}}
+#pragma acc declare copy(NSVar1) copyin(NSVar2), copyout(NSVar3), create(NSVar4), present(NSVar5), deviceptr(NSVar6), device_resident(NSVar7), link(NSVar8)
+
+extern "C" {
+  int ExternVar, ExternVar1, ExternVar2, ExternVar3, ExternVar4, *ExternVar5, ExternVar6, ExternVar7;
+  // Only create, copyin, deviceptr, device-resident, link at NS scope.
+  // expected-error at +3{{OpenACC 'copy' clause on a 'declare' directive is not allowed at global or namespace scope}}
+  // expected-error at +2{{OpenACC 'copyout' clause on a 'declare' directive is not allowed at global or namespace scope}}
+  // expected-error at +1{{OpenACC 'present' clause on a 'declare' directive is not allowed at global or namespace scope}}
+#pragma acc declare copy(ExternVar) copyin(ExternVar1), copyout(ExternVar2), create(ExternVar3), present(ExternVar4), deviceptr(ExternVar5), device_resident(ExternVar6), link(ExternVar7)
+  }
+}
+// expected-error at +2{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+// expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(NS::NSVar, NS::NSArray)
+
+struct Struct {
+  static const int StaticMem = 5;
+  static const int StaticMem2 = 5;
+  int NonStaticMem;
+// expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(Global)
+  // OK, same scope.
+#pragma acc declare create(StaticMem, StaticMem2)
+// expected-error at +4{{variable referenced in 'create' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -2{{previous reference is here}}
+// expected-error at +2{{variable referenced in 'copyin' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -4{{previous reference is here}}
+#pragma acc declare create(StaticMem) copyin(StaticMem)
+// expected-error at +1{{no valid clauses specified in OpenACC 'declare' directive}}
+#pragma acc declare
+
+  void Inline(int Arg) {
+// expected-error at +1{{no valid clauses specified in OpenACC 'declare' directive}}
+#pragma acc declare
+// expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(StaticMem)
+
+    int Local, Local2, Local3, Local4;
+  // OK, same scope.
+#pragma acc declare create(Local, Arg)
+// expected-error at +2{{variable referenced in 'copyin' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at +1{{previous reference is here}}
+#pragma acc declare create(Local2) copyin(Local2)
+
+    for (int I = 0; I < 5; ++I) {
+      int Other;
+    // FIXME: We don't catch this because we use decl-context instead of scope.
+#pragma acc declare create(Local3, Local4)
+      // OK, same scope.
+#pragma acc declare create(I, Other)
+// expected-error at +4 2{{variable referenced in 'create' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -2 2{{previous reference is here}}
+// expected-error at +2 2{{variable referenced in 'copyin' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -4 2{{previous reference is here}}
+#pragma acc declare create(I, Other) copyin(I, Other)
+
+// expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(NonStaticMem)
+    }
+  }
+
+  void OutOfLine(int Arg, int Arg2);
+};
+
+void Struct::OutOfLine(int Arg, int Arg2) {
+// expected-error at +1{{no valid clauses specified in OpenACC 'declare' directive}}
+#pragma acc declare
+// expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(StaticMem)
+
+  int Local, Local2;
+// OK, same scope.
+#pragma acc declare create(Local, Arg)
+// expected-error at +4{{variable referenced in 'create' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -2{{previous reference is here}}
+// expected-error at +2{{variable referenced in 'copyin' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -4{{previous reference is here}}
+#pragma acc declare create(Local) copyin(Local)
+
+  for (int I = 0; I < 5; ++I) {
+    int Other;
+    // FIXME: We don't catch this because we use decl-context instead of scope.
+#pragma acc declare create(Local2, Arg2)
+    // OK, same scope.
+#pragma acc declare create(I, Other)
+// expected-error at +4 2{{variable referenced in 'create' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -2 2{{previous reference is here}}
+// expected-error at +2 2{{variable referenced in 'copyin' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -4 2{{previous reference is here}}
+#pragma acc declare create(I, Other) copyin(I, Other)
+  }
+// expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(NonStaticMem)
+}
+
+template<typename T>
+struct DepStruct {
+  static const T DepStaticMem = 5;
+  static const int StaticMem = 5;
+  int NonStaticMem;
+// expected-error at +1{{no valid clauses specified in OpenACC 'declare' directive}}
+#pragma acc declare
+// expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(Global)
+  // OK, same scope.
+#pragma acc declare create(DepStaticMem)
+  // OK, same scope.
+#pragma acc declare create(StaticMem)
+// expected-error at +4{{variable referenced in 'create' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -2{{previous reference is here}}
+// expected-error at +2{{variable referenced in 'copyin' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -4{{previous reference is here}}
+#pragma acc declare create(StaticMem) copyin(StaticMem)
+// expected-error at +4{{variable referenced in 'create' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -9{{previous reference is here}}
+// expected-error at +2{{variable referenced in 'copyin' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -11{{previous reference is here}}
+#pragma acc declare create(DepStaticMem) copyin(DepStaticMem)
+
+  void Inline(int Arg) {
+// expected-error at +1{{no valid clauses specified in OpenACC 'declare' directive}}
+#pragma acc declare
+// expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(DepStaticMem)
+// expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(StaticMem)
+
+    T Local, Local2;
+  // OK, same scope.
+#pragma acc declare create(Local, Arg)
+// expected-error at +2 2{{variable referenced in 'create' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -2 2{{previous reference is here}}
+#pragma acc declare create(Local, Local)
+
+    for (int I = 0; I < 5; ++I) {
+      int Other;
+      // FIXME: Since we approximate this as a decl-context, we can't check
+      // scope here.
+#pragma acc declare create(Local2)
+      // OK, same scope.
+#pragma acc declare create(I, Other)
+      // expected-error at +2 3{{variable referenced in 'create' clause of OpenACC 'declare' directive was already referenced}}
+      // expected-note at -2 3{{previous reference is here}}
+#pragma acc declare create(I, Other, I)
+    }
+    // expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(NonStaticMem)
+  }
+
+  void OutOfLine(int Arg);
+
+  template<typename U>
+  void TemplInline(U Arg, U Arg2) {
+// expected-error at +1{{no valid clauses specified in OpenACC 'declare' directive}}
+#pragma acc declare
+    // expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(DepStaticMem)
+    // expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(StaticMem)
+
+    T Local, Local2, Local3;
+  // OK, same scope.
+#pragma acc declare create(Local, Arg)
+// expected-error at +4{{variable referenced in 'create' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -2{{previous reference is here}}
+// expected-error at +2{{variable referenced in 'present' clause of OpenACC 'declare' directive was already referenced}}
+// expected-note at -4{{previous reference is here}}
+#pragma acc declare create(Local2, Arg) present(Local, Arg2)
+    {
+      // FIXME: We don't catch this, since we check decl-context not scopes.
+#pragma acc declare create(Local3)
+
+      // expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(NonStaticMem)
+    }
+  }
+  template<typename U>
+  void TemplOutline(U Arg);
+};
+
+template<typename T>
+void DepStruct<T>::OutOfLine(int Arg) {
+  // expected-error at +1{{no valid clauses specified in OpenACC 'declare' directive}}
+#pragma acc declare
+  // expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(StaticMem)
+
+  T Local, Local2;
+// OK, same scope.
+#pragma acc declare create(Local, Arg)
+
+  for (int I = 0; I < 5; ++I) {
+    int Other;
+    // FIXME: We don't catch this because we use decl-context instead of scope.
+#pragma acc declare create(Local2)
+    // OK, same scope.
+#pragma acc declare create(I, Other)
+  }
+  // expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(NonStaticMem)
+}
+template<typename T>
+template<typename U>
+void DepStruct<T>::TemplOutline(U Arg) {
+// expected-error at +1{{no valid clauses specified in OpenACC 'declare' directive}}
+#pragma acc declare
+  // expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(DepStaticMem)
+  // expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(StaticMem)
+
+  T Local, Local2;
+// OK, same scope.
+#pragma acc declare create(Local, Arg)
+
+  {
+    // FIXME: We could potentially fix this someday, but as we don't have
+    // 'scope' information like this during template instantiation, we have to
+    // permit this.
+#pragma acc declare create(Local2)
+  }
+  // expected-error at +1{{variable appearing in 'create' clause of OpenACC 'declare' directive must be in the same scope as the directive}}
+#pragma acc declare create(NonStaticMem)
+}
+
+void use() {
+  DepStruct<int> DS;
+  DS.Inline(1);
+  DS.OutOfLine(1);
+  DS.TemplInline(1, 2);
+  DS.TemplOutline(1);
+}
+
+// Only variable or array name.
+
+// expected-error at +1{{OpenACC variable on 'declare' construct is not a valid variable name or array name}}
+#pragma acc declare create(GlobalArray[0])
+// expected-error at +1{{OpenACC variable on 'declare' construct is not a valid variable name or array name}}
+#pragma acc declare create(GlobalArray[0: 1])
+
+struct S { int I; };
+// expected-error at +1{{OpenACC variable on 'declare' construct is not a valid variable name or array name}}
+#pragma acc declare create(S{}.I)
+
+int GS1, GS2, GS3, GS4, GS5, *GS6, GS7, GS8;
+
+// Only create, copyin, deviceptr, device-resident, link at NS scope.
+// expected-error at +3{{OpenACC 'copy' clause on a 'declare' directive is not allowed at global or namespace scope}}
+// expected-error at +2{{OpenACC 'copyout' clause on a 'declare' directive is not allowed at global or namespace scope}}
+// expected-error at +1{{OpenACC 'present' clause on a 'declare' directive is not allowed at global or namespace scope}}
+#pragma acc declare copy(GS1) copyin(GS2), copyout(GS3), create(GS4), present(GS5), deviceptr(GS6), device_resident(GS7), link(GS8)
+
+void ExternVar() {
+  extern int I, I2, I3, I4, I5, *I6, I7, I8;
+// expected-error at +3{{'extern' variable may not be referenced by 'copy' clause on an OpenACC 'declare' directive}}
+// expected-error at +2{{'extern' variable may not be referenced by 'copyout' clause on an OpenACC 'declare' directive}}
+// expected-error at +1{{'extern' variable may not be referenced by 'present' clause on an OpenACC 'declare' directive}}
+#pragma acc declare copy(I) copyin(I2), copyout(I3), create(I4), present(I5), deviceptr(I6), device_resident(I7), link(I8)
+}
+
+// Link can only have global, namespace, or extern vars.
+#pragma acc declare link(Global, GlobalArray)
+
+struct Struct2 {
+  static const int StaticMem = 5;
+  // expected-error at +1{{variable referenced by 'link' clause not in global or namespace scope must be marked 'extern'}}
+#pragma acc declare link(StaticMem)
+
+  void MemFunc(int I) {
+    int Local;
+    extern int ExternLocal;
+
+  // expected-error at +2{{variable referenced by 'link' clause not in global or namespace scope must be marked 'extern'}}
+  // expected-error at +1{{variable referenced by 'link' clause not in global or namespace scope must be marked 'extern'}}
+#pragma acc declare link(I, Local, ExternLocal)
+}
+};
+

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 f56a1267fbad1..b4a705ba7d1c4 100644
--- a/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
+++ b/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
@@ -86,7 +86,7 @@ void uses() {
   // expected-error at +1{{OpenACC 'deviceptr' clause is not valid on 'loop' directive}}
 #pragma acc loop auto deviceptr(VarPtr)
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented}}
+  // expected-error at +1{{OpenACC 'device_resident' clause is not valid on 'loop' directive}}
 #pragma acc loop auto device_resident(VarPtr)
   for(unsigned i = 0; i < 5; ++i);
   // expected-error at +1{{OpenACC 'firstprivate' clause is not valid on 'loop' directive}}
@@ -95,7 +95,7 @@ void uses() {
   // expected-error at +1{{OpenACC 'host' clause is not valid on 'loop' directive}}
 #pragma acc loop auto host(Var)
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning at +1{{OpenACC clause 'link' not yet implemented}}
+  // expected-error at +1{{OpenACC 'link' clause is not valid on 'loop' directive}}
 #pragma acc loop auto link(Var)
   for(unsigned i = 0; i < 5; ++i);
   // expected-error at +1{{OpenACC 'no_create' clause is not valid on 'loop' directive}}
@@ -220,7 +220,7 @@ void uses() {
   // expected-error at +1{{OpenACC 'deviceptr' clause is not valid on 'loop' directive}}
 #pragma acc loop deviceptr(VarPtr) auto
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented}}
+  // expected-error at +1{{OpenACC 'device_resident' clause is not valid on 'loop' directive}}
 #pragma acc loop device_resident(VarPtr) auto
   for(unsigned i = 0; i < 5; ++i);
   // expected-error at +1{{OpenACC 'firstprivate' clause is not valid on 'loop' directive}}
@@ -229,7 +229,7 @@ void uses() {
   // expected-error at +1{{OpenACC 'host' clause is not valid on 'loop' directive}}
 #pragma acc loop host(Var) auto
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning at +1{{OpenACC clause 'link' not yet implemented}}
+  // expected-error at +1{{OpenACC 'link' clause is not valid on 'loop' directive}}
 #pragma acc loop link(Var) auto
   for(unsigned i = 0; i < 5; ++i);
   // expected-error at +1{{OpenACC 'no_create' clause is not valid on 'loop' directive}}
@@ -355,7 +355,7 @@ void uses() {
   // expected-error at +1{{OpenACC 'deviceptr' clause is not valid on 'loop' directive}}
 #pragma acc loop independent deviceptr(VarPtr)
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented}}
+  // expected-error at +1{{OpenACC 'device_resident' clause is not valid on 'loop' directive}}
 #pragma acc loop independent device_resident(VarPtr)
   for(unsigned i = 0; i < 5; ++i);
   // expected-error at +1{{OpenACC 'firstprivate' clause is not valid on 'loop' directive}}
@@ -364,7 +364,7 @@ void uses() {
   // expected-error at +1{{OpenACC 'host' clause is not valid on 'loop' directive}}
 #pragma acc loop independent host(Var)
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning at +1{{OpenACC clause 'link' not yet implemented}}
+  // expected-error at +1{{OpenACC 'link' clause is not valid on 'loop' directive}}
 #pragma acc loop independent link(Var)
   for(unsigned i = 0; i < 5; ++i);
   // expected-error at +1{{OpenACC 'no_create' clause is not valid on 'loop' directive}}
@@ -489,7 +489,7 @@ void uses() {
   // expected-error at +1{{OpenACC 'deviceptr' clause is not valid on 'loop' directive}}
 #pragma acc loop deviceptr(VarPtr) independent
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented}}
+  // expected-error at +1{{OpenACC 'device_resident' clause is not valid on 'loop' directive}}
 #pragma acc loop device_resident(VarPtr) independent
   for(unsigned i = 0; i < 5; ++i);
   // expected-error at +1{{OpenACC 'firstprivate' clause is not valid on 'loop' directive}}
@@ -498,7 +498,7 @@ void uses() {
   // expected-error at +1{{OpenACC 'host' clause is not valid on 'loop' directive}}
 #pragma acc loop host(Var) independent
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning at +1{{OpenACC clause 'link' not yet implemented}}
+  // expected-error at +1{{OpenACC 'link' clause is not valid on 'loop' directive}}
 #pragma acc loop link(Var) independent
   for(unsigned i = 0; i < 5; ++i);
   // expected-error at +1{{OpenACC 'no_create' clause is not valid on 'loop' directive}}
@@ -632,7 +632,7 @@ void uses() {
   // expected-error at +1{{OpenACC 'deviceptr' clause is not valid on 'loop' directive}}
 #pragma acc loop seq deviceptr(VarPtr)
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented}}
+  // expected-error at +1{{OpenACC 'device_resident' clause is not valid on 'loop' directive}}
 #pragma acc loop seq device_resident(VarPtr)
   for(unsigned i = 0; i < 5; ++i);
   // expected-error at +1{{OpenACC 'firstprivate' clause is not valid on 'loop' directive}}
@@ -641,7 +641,7 @@ void uses() {
   // expected-error at +1{{OpenACC 'host' clause is not valid on 'loop' directive}}
 #pragma acc loop seq host(Var)
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning at +1{{OpenACC clause 'link' not yet implemented}}
+  // expected-error at +1{{OpenACC 'link' clause is not valid on 'loop' directive}}
 #pragma acc loop seq link(Var)
   for(unsigned i = 0; i < 5; ++i);
   // expected-error at +1{{OpenACC 'no_create' clause is not valid on 'loop' directive}}
@@ -772,7 +772,7 @@ void uses() {
   // expected-error at +1{{OpenACC 'deviceptr' clause is not valid on 'loop' directive}}
 #pragma acc loop deviceptr(VarPtr) seq
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning at +1{{OpenACC clause 'device_resident' not yet implemented}}
+  // expected-error at +1{{OpenACC 'device_resident' clause is not valid on 'loop' directive}}
 #pragma acc loop device_resident(VarPtr) seq
   for(unsigned i = 0; i < 5; ++i);
   // expected-error at +1{{OpenACC 'firstprivate' clause is not valid on 'loop' directive}}
@@ -781,7 +781,7 @@ void uses() {
   // expected-error at +1{{OpenACC 'host' clause is not valid on 'loop' directive}}
 #pragma acc loop host(Var) seq
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning at +1{{OpenACC clause 'link' not yet implemented}}
+  // expected-error at +1{{OpenACC 'link' clause is not valid on 'loop' directive}}
 #pragma acc loop link(Var) seq
   for(unsigned i = 0; i < 5; ++i);
   // expected-error at +1{{OpenACC 'no_create' 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 2c1189bc647d0..f05ca11b143c0 100644
--- a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
@@ -98,8 +98,7 @@ void uses() {
   // expected-error at +1{{OpenACC 'deviceptr' clause is not valid on 'loop' directive}}
 #pragma acc loop device_type(*) deviceptr(VarPtr)
   for(int i = 0; i < 5; ++i);
-  // expected-error at +2{{OpenACC clause 'device_resident' may not follow a 'device_type' clause in a 'loop' construct}}
-  // expected-note at +1{{previous clause is here}}
+  // expected-error at +1{{OpenACC 'device_resident' clause is not valid on 'loop' directive}}
 #pragma acc loop device_type(*)  device_resident(VarPtr)
   for(int i = 0; i < 5; ++i);
   // expected-error at +1{{OpenACC 'firstprivate' clause is not valid on 'loop' directive}}
@@ -108,8 +107,7 @@ void uses() {
   // expected-error at +1{{OpenACC 'host' clause is not valid on 'loop' directive}}
 #pragma acc loop device_type(*) host(Var)
   for(int i = 0; i < 5; ++i);
-  // expected-error at +2{{OpenACC clause 'link' may not follow a 'device_type' clause in a 'loop' construct}}
-  // expected-note at +1{{previous clause is here}}
+  // expected-error at +1{{OpenACC 'link' clause is not valid on 'loop' directive}}
 #pragma acc loop device_type(*) link(Var)
   for(int i = 0; i < 5; ++i);
   // expected-error at +1{{OpenACC 'no_create' clause is not valid on 'loop' directive}}

diff  --git a/clang/test/SemaOpenACC/unimplemented-construct.c b/clang/test/SemaOpenACC/unimplemented-construct.c
index b22bfa5a69bb0..261c855cdd010 100644
--- a/clang/test/SemaOpenACC/unimplemented-construct.c
+++ b/clang/test/SemaOpenACC/unimplemented-construct.c
@@ -4,36 +4,36 @@
 #pragma acc routine
 
 struct S {
-// expected-warning at +1{{OpenACC construct 'declare' not yet implemented, pragma ignored}}
-#pragma acc declare 
+// expected-warning at +1{{OpenACC construct 'routine' not yet implemented, pragma ignored}}
+#pragma acc routine
 int foo;
 };
 
 void func() {
-// expected-warning at +1{{OpenACC construct 'declare' not yet implemented, pragma ignored}}
-#pragma acc declare
+// expected-warning at +1{{OpenACC construct 'routine' not yet implemented, pragma ignored}}
+#pragma acc routine
   int foo;
 
-// expected-warning at +1{{OpenACC construct 'declare' not yet implemented, pragma ignored}}
-#pragma acc declare
+// expected-warning at +1{{OpenACC construct 'routine' not yet implemented, pragma ignored}}
+#pragma acc routine
   {
-// expected-warning at +1{{OpenACC construct 'declare' not yet implemented, pragma ignored}}
-#pragma acc declare
+// expected-warning at +1{{OpenACC construct 'routine' not yet implemented, pragma ignored}}
+#pragma acc routine
     {
-// expected-warning at +1{{OpenACC construct 'declare' not yet implemented, pragma ignored}}
-#pragma acc declare
+// expected-warning at +1{{OpenACC construct 'routine' not yet implemented, pragma ignored}}
+#pragma acc routine
     }
   }
 
-// expected-warning at +1{{OpenACC construct 'declare' not yet implemented, pragma ignored}}
-#pragma acc declare
+// expected-warning at +1{{OpenACC construct 'routine' not yet implemented, pragma ignored}}
+#pragma acc routine
   while(0){}
 
-// expected-warning at +1{{OpenACC construct 'declare' not yet implemented, pragma ignored}}
-#pragma acc declare
+// expected-warning at +1{{OpenACC construct 'routine' not yet implemented, pragma ignored}}
+#pragma acc routine
   for(;;){}
 
-// expected-warning at +1{{OpenACC construct 'declare' not yet implemented, pragma ignored}}
-#pragma acc declare
+// expected-warning at +1{{OpenACC construct 'routine' not yet implemented, pragma ignored}}
+#pragma acc routine
 };
 

diff  --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
index c710816bd24d0..479490442f3c8 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -2905,6 +2905,13 @@ void OpenACCClauseEnqueue::VisitNoCreateClause(const OpenACCNoCreateClause &C) {
 void OpenACCClauseEnqueue::VisitCopyClause(const OpenACCCopyClause &C) {
   VisitVarList(C);
 }
+void OpenACCClauseEnqueue::VisitLinkClause(const OpenACCLinkClause &C) {
+  VisitVarList(C);
+}
+void OpenACCClauseEnqueue::VisitDeviceResidentClause(
+    const OpenACCDeviceResidentClause &C) {
+  VisitVarList(C);
+}
 void OpenACCClauseEnqueue::VisitCopyInClause(const OpenACCCopyInClause &C) {
   VisitVarList(C);
 }
@@ -7249,6 +7256,7 @@ CXCursor clang_getCursorDefinition(CXCursor C) {
   case Decl::LifetimeExtendedTemporary:
   case Decl::RequiresExprBody:
   case Decl::UnresolvedUsingIfExists:
+  case Decl::OpenACCDeclare:
     return C;
 
   // Declaration kinds that don't make any sense here, but are


        


More information about the cfe-commits mailing list