diff --git a/clang/include/clang/AST/OpenMPClause.h b/clang/include/clang/AST/OpenMPClause.h
--- a/clang/include/clang/AST/OpenMPClause.h
+++ b/clang/include/clang/AST/OpenMPClause.h
@@ -7416,6 +7416,110 @@
   }
 };
 
+/// This represents clause 'has_device_ptr' in the '#pragma omp ...'
+/// directives.
+///
+/// \code
+/// #pragma omp target has_device_addr(a,b)
+/// \endcode
+/// In this example directive '#pragma omp target' has clause
+/// 'has_device_ptr' with the variables 'a' and 'b'.
+class OMPHasDeviceAddrClause final
+    : public OMPMappableExprListClause<OMPHasDeviceAddrClause>,
+      private llvm::TrailingObjects<
+          OMPHasDeviceAddrClause, Expr *, ValueDecl *, unsigned,
+          OMPClauseMappableExprCommon::MappableComponent> {
+  friend class OMPClauseReader;
+  friend OMPMappableExprListClause;
+  friend OMPVarListClause;
+  friend TrailingObjects;
+
+  /// Build clause with number of variables \a NumVars.
+  ///
+  /// \param Locs Locations needed to build a mappable clause. It includes 1)
+  /// StartLoc: starting location of the clause (the clause keyword); 2)
+  /// LParenLoc: location of '('; 3) EndLoc: ending location of the clause.
+  /// \param Sizes All required sizes to build a mappable clause. It includes 1)
+  /// NumVars: number of expressions listed in this clause; 2)
+  /// NumUniqueDeclarations: number of unique base declarations in this clause;
+  /// 3) NumComponentLists: number of component lists in this clause; and 4)
+  /// NumComponents: total number of expression components in the clause.
+  explicit OMPHasDeviceAddrClause(const OMPVarListLocTy &Locs,
+                                  const OMPMappableExprListSizeTy &Sizes)
+      : OMPMappableExprListClause(llvm::omp::OMPC_has_device_addr, Locs,
+                                  Sizes) {}
+
+  /// Build an empty clause.
+  ///
+  /// \param Sizes All required sizes to build a mappable clause. It includes 1)
+  /// NumVars: number of expressions listed in this clause; 2)
+  /// NumUniqueDeclarations: number of unique base declarations in this clause;
+  /// 3) NumComponentLists: number of component lists in this clause; and 4)
+  /// NumComponents: total number of expression components in the clause.
+  explicit OMPHasDeviceAddrClause(const OMPMappableExprListSizeTy &Sizes)
+      : OMPMappableExprListClause(llvm::omp::OMPC_has_device_addr,
+                                  OMPVarListLocTy(), Sizes) {}
+
+  /// Define the sizes of each trailing object array except the last one. This
+  /// is required for TrailingObjects to work properly.
+  size_t numTrailingObjects(OverloadToken<Expr *>) const {
+    return varlist_size();
+  }
+  size_t numTrailingObjects(OverloadToken<ValueDecl *>) const {
+    return getUniqueDeclarationsNum();
+  }
+  size_t numTrailingObjects(OverloadToken<unsigned>) const {
+    return getUniqueDeclarationsNum() + getTotalComponentListNum();
+  }
+
+public:
+  /// Creates clause with a list of variables \a Vars.
+  ///
+  /// \param C AST context.
+  /// \param Locs Locations needed to build a mappable clause. It includes 1)
+  /// StartLoc: starting location of the clause (the clause keyword); 2)
+  /// LParenLoc: location of '('; 3) EndLoc: ending location of the clause.
+  /// \param Vars The original expression used in the clause.
+  /// \param Declarations Declarations used in the clause.
+  /// \param ComponentLists Component lists used in the clause.
+  static OMPHasDeviceAddrClause *
+  Create(const ASTContext &C, const OMPVarListLocTy &Locs,
+         ArrayRef<Expr *> Vars, ArrayRef<ValueDecl *> Declarations,
+         MappableExprComponentListsRef ComponentLists);
+
+  /// Creates an empty clause with the place for \a NumVars variables.
+  ///
+  /// \param C AST context.
+  /// \param Sizes All required sizes to build a mappable clause. It includes 1)
+  /// NumVars: number of expressions listed in this clause; 2)
+  /// NumUniqueDeclarations: number of unique base declarations in this clause;
+  /// 3) NumComponentLists: number of component lists in this clause; and 4)
+  /// NumComponents: total number of expression components in the clause.
+  static OMPHasDeviceAddrClause *
+  CreateEmpty(const ASTContext &C, const OMPMappableExprListSizeTy &Sizes);
+
+  child_range children() {
+    return child_range(reinterpret_cast<Stmt **>(varlist_begin()),
+                       reinterpret_cast<Stmt **>(varlist_end()));
+  }
+
+  const_child_range children() const {
+    auto Children = const_cast<OMPHasDeviceAddrClause *>(this)->children();
+    return const_child_range(Children.begin(), Children.end());
+  }
+
+  child_range used_children() {
+    return child_range(child_iterator(), child_iterator());
+  }
+  const_child_range used_children() const {
+    return const_child_range(const_child_iterator(), const_child_iterator());
+  }
+
+  static bool classof(const OMPClause *T) {
+    return T->getClauseKind() == llvm::omp::OMPC_has_device_addr;
+  }
+};
+
 /// This represents clause 'nontemporal' in the '#pragma omp ...' directives.
 ///
 /// \code
diff --git a/clang/include/clang/AST/RecursiveASTVisitor.h b/clang/include/clang/AST/RecursiveASTVisitor.h
--- a/clang/include/clang/AST/RecursiveASTVisitor.h
+++ b/clang/include/clang/AST/RecursiveASTVisitor.h
@@ -3702,6 +3702,13 @@
   return true;
 }
 
+template <typename Derived>
+bool RecursiveASTVisitor<Derived>::VisitOMPHasDeviceAddrClause(
+    OMPHasDeviceAddrClause *C) {
+  TRY_TO(VisitOMPClauseList(C));
+  return true;
+}
+
 template <typename Derived>
 bool RecursiveASTVisitor<Derived>::VisitOMPNontemporalClause(
     OMPNontemporalClause *C) {
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -11541,6 +11541,9 @@
   /// Called on well-formed 'is_device_ptr' clause.
   OMPClause *ActOnOpenMPIsDevicePtrClause(ArrayRef<Expr *> VarList,
                                           const OMPVarListLocTy &Locs);
+  /// Called on well-formed 'has_device_addr' clause.
+  OMPClause *ActOnOpenMPHasDeviceAddrClause(ArrayRef<Expr *> VarList,
+                                            const OMPVarListLocTy &Locs);
   /// Called on well-formed 'nontemporal' clause.
   OMPClause *ActOnOpenMPNontemporalClause(ArrayRef<Expr *> VarList,
                                           SourceLocation StartLoc,
diff --git a/clang/lib/AST/OpenMPClause.cpp b/clang/lib/AST/OpenMPClause.cpp
--- a/clang/lib/AST/OpenMPClause.cpp
+++ b/clang/lib/AST/OpenMPClause.cpp
@@ -146,6 +146,7 @@
   case OMPC_use_device_ptr:
   case OMPC_use_device_addr:
   case OMPC_is_device_ptr:
+  case OMPC_has_device_addr:
   case OMPC_unified_address:
   case OMPC_unified_shared_memory:
   case OMPC_reverse_offload:
@@ -244,6 +245,7 @@
   case OMPC_use_device_ptr:
   case OMPC_use_device_addr:
   case OMPC_is_device_ptr:
+  case OMPC_has_device_addr:
   case OMPC_unified_address:
   case OMPC_unified_shared_memory:
   case OMPC_reverse_offload:
@@ -1432,6 +1434,53 @@
   return new (Mem) OMPIsDevicePtrClause(Sizes);
 }
 
+OMPHasDeviceAddrClause *
+OMPHasDeviceAddrClause::Create(const ASTContext &C, const OMPVarListLocTy &Locs,
+                               ArrayRef<Expr *> Vars,
+                               ArrayRef<ValueDecl *> Declarations,
+                               MappableExprComponentListsRef ComponentLists) {
+  OMPMappableExprListSizeTy Sizes;
+  Sizes.NumVars = Vars.size();
+  Sizes.NumUniqueDeclarations = getUniqueDeclarationsTotalNumber(Declarations);
+  Sizes.NumComponentLists = ComponentLists.size();
+  Sizes.NumComponents = getComponentsTotalNumber(ComponentLists);
+
+  // We need to allocate:
+  // NumVars x Expr* - we have an original list expression for each clause list
+  // entry.
+  // NumUniqueDeclarations x ValueDecl* - unique base declarations associated
+  // with each component list.
+  // (NumUniqueDeclarations + NumComponentLists) x unsigned - we specify the
+  // number of lists for each unique declaration and the size of each component
+  // list.
+  // NumComponents x MappableComponent - the total of all the components in all
+  // the lists.
+  void *Mem = C.Allocate(
+      totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
+                       OMPClauseMappableExprCommon::MappableComponent>(
+          Sizes.NumVars, Sizes.NumUniqueDeclarations,
+          Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
+          Sizes.NumComponents));
+
+  auto *Clause = new (Mem) OMPHasDeviceAddrClause(Locs, Sizes);
+
+  Clause->setVarRefs(Vars);
+  Clause->setClauseInfo(Declarations, ComponentLists);
+  return Clause;
+}
+
+OMPHasDeviceAddrClause *
+OMPHasDeviceAddrClause::CreateEmpty(const ASTContext &C,
+                                    const OMPMappableExprListSizeTy &Sizes) {
+  void *Mem = C.Allocate(
+      totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
+                       OMPClauseMappableExprCommon::MappableComponent>(
+          Sizes.NumVars, Sizes.NumUniqueDeclarations,
+          Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
+          Sizes.NumComponents));
+  return new (Mem) OMPHasDeviceAddrClause(Sizes);
+}
+
 OMPNontemporalClause *OMPNontemporalClause::Create(const ASTContext &C,
                                                    SourceLocation StartLoc,
                                                    SourceLocation LParenLoc,
@@ -2259,6 +2308,14 @@
   }
 }
 
+void OMPClausePrinter::VisitOMPHasDeviceAddrClause(OMPHasDeviceAddrClause *Node) {
+  if (!Node->varlist_empty()) {
+    OS << "has_device_addr";
+    VisitOMPClauseList(Node, '(');
+    OS << ")";
+  }
+}
+
 void OMPClausePrinter::VisitOMPNontemporalClause(OMPNontemporalClause *Node) {
   if (!Node->varlist_empty()) {
     OS << "nontemporal";
diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -861,6 +861,10 @@
     const OMPIsDevicePtrClause *C) {
   VisitOMPClauseList(C);
 }
+void OMPClauseProfiler::VisitOMPHasDeviceAddrClause(
+    const OMPHasDeviceAddrClause *C) {
+  VisitOMPClauseList(C);
+}
 void OMPClauseProfiler::VisitOMPNontemporalClause(
     const OMPNontemporalClause *C) {
   VisitOMPClauseList(C);
diff --git a/clang/lib/Basic/OpenMPKinds.cpp b/clang/lib/Basic/OpenMPKinds.cpp
--- a/clang/lib/Basic/OpenMPKinds.cpp
+++ b/clang/lib/Basic/OpenMPKinds.cpp
@@ -186,6 +186,7 @@
   case OMPC_use_device_ptr:
   case OMPC_use_device_addr:
   case OMPC_is_device_ptr:
+  case OMPC_has_device_addr:
   case OMPC_unified_address:
   case OMPC_unified_shared_memory:
   case OMPC_reverse_offload:
@@ -452,6 +453,7 @@
   case OMPC_use_device_ptr:
   case OMPC_use_device_addr:
   case OMPC_is_device_ptr:
+  case OMPC_has_device_addr:
   case OMPC_unified_address:
   case OMPC_unified_shared_memory:
   case OMPC_reverse_offload:
diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp
--- a/clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -6164,6 +6164,7 @@
   case OMPC_use_device_ptr:
   case OMPC_use_device_addr:
   case OMPC_is_device_ptr:
+  case OMPC_has_device_addr:
   case OMPC_unified_address:
   case OMPC_unified_shared_memory:
   case OMPC_reverse_offload:
diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp
--- a/clang/lib/Parse/ParseOpenMP.cpp
+++ b/clang/lib/Parse/ParseOpenMP.cpp
@@ -3108,7 +3108,8 @@
 ///       in_reduction-clause | allocator-clause | allocate-clause |
 ///       acq_rel-clause | acquire-clause | release-clause | relaxed-clause |
 ///       depobj-clause | destroy-clause | detach-clause | inclusive-clause |
-///       exclusive-clause | uses_allocators-clause | use_device_addr-clause
+///       exclusive-clause | uses_allocators-clause | use_device_addr-clause |
+///       has_device_addr
 ///
 OMPClause *Parser::ParseOpenMPClause(OpenMPDirectiveKind DKind,
                                      OpenMPClauseKind CKind, bool FirstClause) {
@@ -3290,6 +3291,7 @@
   case OMPC_use_device_ptr:
   case OMPC_use_device_addr:
   case OMPC_is_device_ptr:
+  case OMPC_has_device_addr:
   case OMPC_allocate:
   case OMPC_nontemporal:
   case OMPC_inclusive:
@@ -4449,6 +4451,8 @@
 ///       'use_device_addr' '(' list ')'
 ///    is_device_ptr-clause:
 ///       'is_device_ptr' '(' list ')'
+///    has_device_addr-clause:
+///       'has_device_addr' '(' list ')'
 ///    allocate-clause:
 ///       'allocate' '(' [ allocator ':' ] list ')'
 ///    nontemporal-clause:
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -6431,6 +6431,7 @@
       case OMPC_use_device_ptr:
       case OMPC_use_device_addr:
       case OMPC_is_device_ptr:
+      case OMPC_has_device_addr:
       case OMPC_nontemporal:
       case OMPC_order:
       case OMPC_destroy:
@@ -15959,6 +15960,7 @@
   case OMPC_use_device_ptr:
   case OMPC_use_device_addr:
   case OMPC_is_device_ptr:
+  case OMPC_has_device_addr:
   case OMPC_unified_address:
   case OMPC_unified_shared_memory:
   case OMPC_reverse_offload:
@@ -16264,6 +16266,7 @@
   case OMPC_use_device_ptr:
   case OMPC_use_device_addr:
   case OMPC_is_device_ptr:
+  case OMPC_has_device_addr:
   case OMPC_unified_address:
   case OMPC_unified_shared_memory:
   case OMPC_reverse_offload:
@@ -16523,6 +16526,7 @@
   case OMPC_use_device_ptr:
   case OMPC_use_device_addr:
   case OMPC_is_device_ptr:
+  case OMPC_has_device_addr:
   case OMPC_atomic_default_mem_order:
   case OMPC_device_type:
   case OMPC_match:
@@ -17014,6 +17018,9 @@
   case OMPC_is_device_ptr:
     Res = ActOnOpenMPIsDevicePtrClause(VarList, Locs);
     break;
+  case OMPC_has_device_addr:
+    Res = ActOnOpenMPHasDeviceAddrClause(VarList, Locs);
+    break;
   case OMPC_allocate:
     Res = ActOnOpenMPAllocateClause(DepModOrTailExpr, VarList, StartLoc,
                                     LParenLoc, ColonLoc, EndLoc);
@@ -22426,6 +22433,88 @@
                                       MVLI.VarComponents);
 }
 
+OMPClause *Sema::ActOnOpenMPHasDeviceAddrClause(ArrayRef<Expr *> VarList,
+                                                const OMPVarListLocTy &Locs) {
+  MappableVarListInfo MVLI(VarList);
+  for (Expr *RefExpr : VarList) {
+    assert(RefExpr && "NULL expr in OpenMP has_device_addr clause.");
+    SourceLocation ELoc;
+    SourceRange ERange;
+    Expr *SimpleRefExpr = RefExpr;
+    auto Res = getPrivateItem(*this, SimpleRefExpr, ELoc, ERange,
+                              /*AllowArraySection=*/true);
+    if (Res.second) {
+      // It will be analyzed later.
+      MVLI.ProcessedVarList.push_back(RefExpr);
+    }
+    ValueDecl *D = Res.first;
+    if (!D)
+      continue;
+
+    // Check if the declaration in the clause does not show up in any data
+    // sharing attribute.
+    DSAStackTy::DSAVarData DVar = DSAStack->getTopDSA(D, /*FromParent=*/false);
+    if (isOpenMPPrivate(DVar.CKind)) {
+      Diag(ELoc, diag::err_omp_variable_in_given_clause_and_dsa)
+          << getOpenMPClauseName(DVar.CKind)
+          << getOpenMPClauseName(OMPC_has_device_addr)
+          << getOpenMPDirectiveName(DSAStack->getCurrentDirective());
+      reportOriginalDsa(*this, DSAStack, D, DVar);
+      continue;
+    }
+
+    const Expr *ConflictExpr;
+    if (DSAStack->checkMappableExprComponentListsForDecl(
+            D, /*CurrentRegionOnly=*/true,
+            [&ConflictExpr](
+                OMPClauseMappableExprCommon::MappableExprComponentListRef R,
+                OpenMPClauseKind) -> bool {
+              ConflictExpr = R.front().getAssociatedExpression();
+              return true;
+            })) {
+      Diag(ELoc, diag::err_omp_map_shared_storage) << RefExpr->getSourceRange();
+      Diag(ConflictExpr->getExprLoc(), diag::note_used_here)
+          << ConflictExpr->getSourceRange();
+      continue;
+    }
+
+    // Store the components in the stack so that they can be used to check
+    // against other clauses later on.
+    OMPClauseMappableExprCommon::MappableComponent MC(
+        SimpleRefExpr, D, /*IsNonContiguous=*/false);
+    DSAStack->addMappableExpressionComponents(
+        D, MC, /*WhereFoundClauseKind=*/OMPC_has_device_addr);
+
+    // Record the expression we've just processed.
+    auto *VD = dyn_cast<VarDecl>(D);
+    if (!VD && !CurContext->isDependentContext()) {
+      DeclRefExpr *Ref =
+          buildCapture(*this, D, SimpleRefExpr, /*WithInit=*/true);
+      assert(Ref && "has_device_addr capture failed");
+      MVLI.ProcessedVarList.push_back(Ref);
+    } else
+      MVLI.ProcessedVarList.push_back(RefExpr->IgnoreParens());
+
+    // Create a mappable component for the list item. List items in this clause
+    // only need a component. We use a null declaration to signal fields in
+    // 'this'.
+    assert((isa<DeclRefExpr>(SimpleRefExpr) ||
+            isa<CXXThisExpr>(cast<MemberExpr>(SimpleRefExpr)->getBase())) &&
+           "Unexpected device pointer expression!");
+    MVLI.VarBaseDeclarations.push_back(
+        isa<DeclRefExpr>(SimpleRefExpr) ? D : nullptr);
+    MVLI.VarComponents.resize(MVLI.VarComponents.size() + 1);
+    MVLI.VarComponents.back().push_back(MC);
+  }
+
+  if (MVLI.ProcessedVarList.empty())
+    return nullptr;
+
+  return OMPHasDeviceAddrClause::Create(Context, Locs, MVLI.ProcessedVarList,
+                                        MVLI.VarBaseDeclarations,
+                                        MVLI.VarComponents);
+}
+
 OMPClause *Sema::ActOnOpenMPAllocateClause(
     Expr *Allocator, ArrayRef<Expr *> VarList, SourceLocation StartLoc,
     SourceLocation ColonLoc, SourceLocation LParenLoc, SourceLocation EndLoc) {
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -2101,6 +2101,15 @@
     return getSema().ActOnOpenMPIsDevicePtrClause(VarList, Locs);
   }
 
+  /// Build a new OpenMP 'has_device_addr' clause.
+  ///
+  /// By default, performs semantic analysis to build the new OpenMP clause.
+  /// Subclasses may override this routine to provide different behavior.
+  OMPClause *RebuildOMPHasDeviceAddrClause(ArrayRef<Expr *> VarList,
+                                           const OMPVarListLocTy &Locs) {
+    return getSema().ActOnOpenMPHasDeviceAddrClause(VarList, Locs);
+  }
+
   /// Build a new OpenMP 'defaultmap' clause.
   ///
   /// By default, performs semantic analysis to build the new OpenMP clause.
@@ -10270,6 +10279,21 @@
   return getDerived().RebuildOMPIsDevicePtrClause(Vars, Locs);
 }
 
+template <typename Derived>
+OMPClause *TreeTransform<Derived>::TransformOMPHasDeviceAddrClause(
+    OMPHasDeviceAddrClause *C) {
+  llvm::SmallVector<Expr *, 16> Vars;
+  Vars.reserve(C->varlist_size());
+  for (auto *VE : C->varlists()) {
+    ExprResult EVar = getDerived().TransformExpr(cast<Expr>(VE));
+    if (EVar.isInvalid())
+      return nullptr;
+    Vars.push_back(EVar.get());
+  }
+  OMPVarListLocTy Locs(C->getBeginLoc(), C->getLParenLoc(), C->getEndLoc());
+  return getDerived().RebuildOMPHasDeviceAddrClause(Vars, Locs);
+}
+
 template <typename Derived>
 OMPClause *
 TreeTransform<Derived>::TransformOMPNontemporalClause(OMPNontemporalClause *C) {
diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -11867,6 +11867,15 @@
     C = OMPIsDevicePtrClause::CreateEmpty(Context, Sizes);
     break;
   }
+  case llvm::omp::OMPC_has_device_addr: {
+    OMPMappableExprListSizeTy Sizes;
+    Sizes.NumVars = Record.readInt();
+    Sizes.NumUniqueDeclarations = Record.readInt();
+    Sizes.NumComponentLists = Record.readInt();
+    Sizes.NumComponents = Record.readInt();
+    C = OMPHasDeviceAddrClause::CreateEmpty(Context, Sizes);
+    break;
+  }
   case llvm::omp::OMPC_allocate:
     C = OMPAllocateClause::CreateEmpty(Context, Record.readInt());
     break;
@@ -12827,6 +12836,49 @@
   C->setComponents(Components, ListSizes);
 }
 
+void OMPClauseReader::VisitOMPHasDeviceAddrClause(OMPHasDeviceAddrClause *C) {
+  C->setLParenLoc(Record.readSourceLocation());
+  auto NumVars = C->varlist_size();
+  auto UniqueDecls = C->getUniqueDeclarationsNum();
+  auto TotalLists = C->getTotalComponentListNum();
+  auto TotalComponents = C->getTotalComponentsNum();
+
+  SmallVector<Expr *, 16> Vars;
+  Vars.reserve(NumVars);
+  for (unsigned I = 0; I != NumVars; ++I)
+    Vars.push_back(Record.readSubExpr());
+  C->setVarRefs(Vars);
+  Vars.clear();
+
+  SmallVector<ValueDecl *, 16> Decls;
+  Decls.reserve(UniqueDecls);
+  for (unsigned I = 0; I < UniqueDecls; ++I)
+    Decls.push_back(Record.readDeclAs<ValueDecl>());
+  C->setUniqueDecls(Decls);
+
+  SmallVector<unsigned, 16> ListsPerDecl;
+  ListsPerDecl.reserve(UniqueDecls);
+  for (unsigned I = 0; I < UniqueDecls; ++I)
+    ListsPerDecl.push_back(Record.readInt());
+  C->setDeclNumLists(ListsPerDecl);
+
+  SmallVector<unsigned, 32> ListSizes;
+  ListSizes.reserve(TotalLists);
+  for (unsigned i = 0; i < TotalLists; ++i)
+    ListSizes.push_back(Record.readInt());
+  C->setComponentListSizes(ListSizes);
+
+  SmallVector<OMPClauseMappableExprCommon::MappableComponent, 32> Components;
+  Components.reserve(TotalComponents);
+  for (unsigned I = 0; I < TotalComponents; ++I) {
+    Expr *AssociatedExpr = Record.readSubExpr();
+    auto *AssociatedDecl = Record.readDeclAs<ValueDecl>();
+    Components.emplace_back(AssociatedExpr, AssociatedDecl,
+                            /*IsNonContiguous=*/false);
+  }
+  C->setComponents(Components, ListSizes);
+}
+
 void OMPClauseReader::VisitOMPNontemporalClause(OMPNontemporalClause *C) {
   C->setLParenLoc(Record.readSourceLocation());
   unsigned NumVars = C->varlist_size();
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -6831,6 +6831,26 @@
   }
 }
 
+void OMPClauseWriter::VisitOMPHasDeviceAddrClause(OMPHasDeviceAddrClause *C) {
+  Record.push_back(C->varlist_size());
+  Record.push_back(C->getUniqueDeclarationsNum());
+  Record.push_back(C->getTotalComponentListNum());
+  Record.push_back(C->getTotalComponentsNum());
+  Record.AddSourceLocation(C->getLParenLoc());
+  for (auto *E : C->varlists())
+    Record.AddStmt(E);
+  for (auto *D : C->all_decls())
+    Record.AddDeclRef(D);
+  for (auto N : C->all_num_lists())
+    Record.push_back(N);
+  for (auto N : C->all_lists_sizes())
+    Record.push_back(N);
+  for (auto &M : C->all_components()) {
+    Record.AddStmt(M.getAssociatedExpression());
+    Record.AddDeclRef(M.getAssociatedDeclaration());
+  }
+}
+
 void OMPClauseWriter::VisitOMPUnifiedAddressClause(OMPUnifiedAddressClause *) {}
 
 void OMPClauseWriter::VisitOMPUnifiedSharedMemoryClause(
diff --git a/clang/test/Analysis/cfg-openmp.cpp b/clang/test/Analysis/cfg-openmp.cpp
--- a/clang/test/Analysis/cfg-openmp.cpp
+++ b/clang/test/Analysis/cfg-openmp.cpp
@@ -1,5 +1,126 @@
 // RUN: %clang_analyze_cc1 -analyzer-checker=debug.DumpCFG %s 2>&1 -fopenmp -fopenmp-version=45 | FileCheck %s
 
+// RUN: %clang_analyze_cc1 -analyzer-checker=debug.DumpCFG %s 2>&1 -fopenmp -fopenmp-version=51 | FileCheck %s --check-prefix=OMP51
+
+#if _OPENMP == 202011
+
+// OMP51-LABEL:  void target_has_device_addr(int argc)
+void target_has_device_addr(int argc) {
+// OMP51:   [B1]
+// OMP51-NEXT:   [[#TTD:]]: 5
+// OMP51-NEXT:   [[#TTD+1]]: int x = 5;
+// OMP51-NEXT:   [[#TTD+2]]: x
+// OMP51-NEXT:   [[#TTD+3]]: [B1.[[#TTD+2]]] (ImplicitCastExpr, LValueToRValue, int)
+// OMP51-NEXT:   [[#TTD+4]]: [B1.[[#TTD+6]]]
+// OMP51-NEXT:   [[#TTD+5]]: [B1.[[#TTD+6]]] = [B1.[[#TTD+3]]]
+// OMP51-NEXT:   [[#TTD+6]]: argc
+// OMP51-NEXT:   [[#TTD+7]]: #pragma omp target has_device_addr(x)
+// OMP51-NEXT:   [B1.[[#TTD+5]]]
+  int x = 5;
+#pragma omp target has_device_addr(x)
+   argc = x;
+}
+// OMP51-LABEL: void target_s_has_device_addr(int argc)
+void target_s_has_device_addr(int argc) {
+  int x, cond, fp, rd, lin, step, map;
+// OMP51-DAG:  [B3]
+// OMP51-DAG: [[#TSB:]]: x
+// OMP51-DAG: [[#TSB+1]]: [B3.[[#TSB]]] (ImplicitCastExpr, LValueToRValue, int)
+// OMP51-DAG: [[#TSB+2]]: argc
+// OMP51-DAG: [[#TSB+3]]: [B3.[[#TSB+2]]] = [B3.[[#TSB+1]]]
+// OMP51-DAG:  [B1]
+// OMP51-DAG: [[#TS:]]: cond
+// OMP51-DAG: [[#TS+1]]: [B1.[[#TS]]] (ImplicitCastExpr, LValueToRValue, int)
+// OMP51-DAG: [[#TS+2]]: [B1.[[#TS+1]]] (ImplicitCastExpr, IntegralToBoolean, _Bool)
+// OMP51-DAG: [[#TS+3]]: fp
+// OMP51-DAG: [[#TS+4]]: rd
+// OMP51-DAG: [[#TS+5]]: lin
+// OMP51-DAG: [[#TS+6]]: step
+// OMP51-DAG: [[#TS+7]]: [B1.[[#TS+6]]] (ImplicitCastExpr, LValueToRValue, int)
+// OMP51-DAG: [[#TS+8]]: [B3.[[#TSB+2]]]
+// OMP51-DAG: [[#TS+9]]: [B3.[[#TSB]]]
+// OMP51-DAG: [[#TS+10]]: #pragma omp target simd if(cond) firstprivate(fp) reduction(+: rd) linear(lin: step) has_device_addr(map)
+// OMP51-DAG:    for (int i = 0;
+// OMP51-DAG: [B3.[[#TSB+3]]];
+#pragma omp target simd if(cond) firstprivate(fp) reduction(+:rd) linear(lin: step) has_device_addr(map)
+  for (int i = 0; i < 10; ++i)
+    argc = x;
+}
+// OMP51-LABEL: void target_t_l_has_device_addr(int argc)
+void target_t_l_has_device_addr(int argc) {
+int x, cond, fp, rd, map;
+// OMP51-DAG: [B3]
+// OMP51-DAG: [[#TTDB:]]: x
+// OMP51-DAG: [[#TTDB+1]]: [B3.[[#TTDB]]] (ImplicitCastExpr, LValueToRValue, int)
+// OMP51-DAG: [[#TTDB+2]]: argc
+// OMP51-DAG: [[#TTDB+3]]: [B3.[[#TTDB+2]]] = [B3.[[#TTDB+1]]]
+// OMP51-DAG: [B1]
+// OMP51-DAG: [[#TTD:]]: cond
+// OMP51-DAG: [[#TTD+1]]: [B1.[[#TTD]]] (ImplicitCastExpr, LValueToRValue, int)
+// OMP51-DAG: [[#TTD+2]]: [B1.[[#TTD+1]]] (ImplicitCastExpr, IntegralToBoolean, _Bool)
+// OMP51-DAG: [[#TTD+3]]: fp
+// OMP51-DAG: [[#TTD+4]]: rd
+// OMP51-DAG: [[#TTD+5]]: [B3.[[#TTDB+2]]]
+// OMP51-DAG: [[#TTD+6]]: [B3.[[#TTDB]]]
+// OMP51-DAG: [[#TTD+7]]:  #pragma omp target teams loop if(cond) firstprivate(fp) reduction(+: rd) has_device_addr(map)
+// OMP51-DAG: for (int i = 0;
+// OMP51-DAG:   [B3.[[#TTDB+3]]];
+#pragma omp target teams loop if(cond) firstprivate(fp) reduction(+:rd) has_device_addr(map)
+   for (int i = 0; i <10; ++i)
+     argc = x;
+}
+// OMP51-LABEL:  void target_p_l_has_device_addr(int argc)
+void target_p_l_has_device_addr(int argc) {
+int x, cond, fp, rd, map;
+#pragma omp target parallel loop if(cond) firstprivate(fp) reduction(+:rd) has_device_addr(map)
+// OMP51-DAG: [B3]
+// OMP51-DAG: [[#TTDB:]]: x
+// OMP51-DAG: [[#TTDB+1]]: [B3.[[#TTDB]]] (ImplicitCastExpr, LValueToRValue, int)
+// OMP51-DAG: [[#TTDB+2]]: argc
+// OMP51-DAG: [[#TTDB+3]]: [B3.[[#TTDB+2]]] = [B3.[[#TTDB+1]]]
+// OMP51-DAG: [B1]
+// OMP51-DAG: [[#TTD:]]: cond
+// OMP51-DAG: [[#TTD+1]]: [B1.[[#TTD]]] (ImplicitCastExpr, LValueToRValue, int)
+// OMP51-DAG: [[#TTD+2]]: [B1.[[#TTD+1]]] (ImplicitCastExpr, IntegralToBoolean, _Bool)
+// OMP51-DAG: [[#TTD+3]]: fp
+// OMP51-DAG: [[#TTD+4]]: rd
+// OMP51-DAG: [[#TTD+5]]: [B3.[[#TTDB+2]]]
+// OMP51-DAG: [[#TTD+6]]: [B3.[[#TTDB]]]
+// OMP51-DAG: [[#TTD+7]]: #pragma omp target parallel loop if(cond) firstprivate(fp) reduction(+: rd) has_device_addr(map)
+// OMP51-DAG: for (int i = 0;
+// OMP51-DAG:   [B3.[[#TTDB+3]]];
+  for (int i = 0; i < 10; ++i)
+    argc = x;
+}
+struct SomeKernel {
+  int targetDev;
+  float devPtr;
+  SomeKernel();
+  ~SomeKernel();
+// OMP51-LABEL: template<> void apply<32U>()
+  template<unsigned int nRHS>
+  void apply() {
+// OMP51-DAG: [B1]
+// OMP51-DAG: [[#TTD:]]: 10
+// OMP51-DAG: [[#TTD+1]]: [B1.[[#TTD:]]] (ImplicitCastExpr, IntegralToFloating, float)
+// OMP51-DAG: [[#TTD+2]]: this
+// OMP51-DAG: [[#TTD+3]]: [B1.[[#TTD+2]]]->devPtr
+// OMP51-DAG: [[#TTD+4]]: [B1.[[#TTD+3]]] = [B1.[[#TTD+1]]]
+// OMP51-DAG: [[#TTD+5]]: #pragma omp target has_device_addr(this->devPtr) device(this->targetDev)
+// OMP51-DAG:    {
+// OMP51-DAG:    [B1.[[#TTD+4]]];
+    #pragma omp target has_device_addr(devPtr) device(targetDev)
+    {
+      devPtr = 10;
+    }
+  }
+};
+void use_template() {
+  SomeKernel aKern;
+  aKern.apply<32>();
+}
+#else // _OPENMP
+
 // CHECK-LABEL:  void xxx(int argc)
 void xxx(int argc) {
 // CHECK:        [B1]
@@ -771,3 +892,5 @@
   for (int i = 0; i < 10; ++i)
     argc = x;
 }
+
+#endif  // _OPENMP
diff --git a/clang/test/OpenMP/target_has_device_addr_ast_print.cpp b/clang/test/OpenMP/target_has_device_addr_ast_print.cpp
new file mode 100644
--- /dev/null
+++ b/clang/test/OpenMP/target_has_device_addr_ast_print.cpp
@@ -0,0 +1,338 @@
+// RUN: %clang_cc1 -verify -fopenmp -std=c++11 -fopenmp-version=51 \
+// RUN:  -ast-print %s | FileCheck %s
+
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -fopenmp-version=51 \
+// RUN:  -emit-pch -o %t %s
+
+// RUN: %clang_cc1 -fopenmp -std=c++11 -fopenmp-version=51 \
+// RUN:  -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+
+// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=51 \
+// RUN:  -std=c++11 -ast-print %s | FileCheck %s
+
+// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 \
+// RUN:  -fopenmp-version=51 -emit-pch -o %t %s
+
+// RUN: %clang_cc1 -fopenmp-simd -std=c++11 -fopenmp-version=51 \
+// RUN:  -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+struct ST {
+  int *a;
+};
+typedef int arr[10];
+typedef ST STarr[10];
+struct SA {
+  const int da[5] = { 0 };
+  ST g[10];
+  STarr &rg = g;
+  int i;
+  int &j = i;
+  int *k = &j;
+  int *&z = k;
+  int aa[10];
+  arr &raa = aa;
+  void func(int arg) {
+#pragma omp target has_device_addr(k)
+    {}
+#pragma omp target has_device_addr(z)
+    {}
+#pragma omp target has_device_addr(aa) // OK
+    {}
+#pragma omp target has_device_addr(raa) // OK
+    {}
+#pragma omp target has_device_addr(g) // OK
+    {}
+#pragma omp target has_device_addr(rg) // OK
+    {}
+#pragma omp target has_device_addr(da) // OK
+    {}
+  return;
+ }
+};
+// CHECK: struct SA
+// CHECK-NEXT: const int da[5] = {0};
+// CHECK-NEXT: ST g[10];
+// CHECK-NEXT: STarr &rg = this->g;
+// CHECK-NEXT: int i;
+// CHECK-NEXT: int &j = this->i;
+// CHECK-NEXT: int *k = &this->j;
+// CHECK-NEXT: int *&z = this->k;
+// CHECK-NEXT: int aa[10];
+// CHECK-NEXT: arr &raa = this->aa;
+// CHECK-NEXT: func(
+// CHECK-NEXT: #pragma omp target has_device_addr(this->k)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(this->z)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(this->aa)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(this->raa)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(this->g)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(this->rg)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(this->da)
+
+struct SB {
+  unsigned A;
+  unsigned B;
+  float Arr[100];
+  float *Ptr;
+  float *foo() {
+    return &Arr[0];
+  }
+};
+
+struct SC {
+  unsigned A : 2;
+  unsigned B : 3;
+  unsigned C;
+  unsigned D;
+  float Arr[100];
+  SB S;
+  SB ArrS[100];
+  SB *PtrS;
+  SB *&RPtrS;
+  float *Ptr;
+
+  SC(SB *&_RPtrS) : RPtrS(_RPtrS) {}
+};
+
+union SD {
+  unsigned A;
+  float B;
+};
+
+struct S1;
+extern S1 a;
+class S2 {
+  mutable int a;
+public:
+  S2():a(0) { }
+  S2(S2 &s2):a(s2.a) { }
+  static float S2s;
+  static const float S2sc;
+};
+const float S2::S2sc = 0;
+const S2 b;
+const S2 ba[5];
+class S3 {
+  int a;
+public:
+  S3():a(0) { }
+  S3(S3 &s3):a(s3.a) { }
+};
+const S3 c;
+const S3 ca[5];
+extern const int f;
+class S4 {
+  int a;
+  S4();
+  S4(const S4 &s4);
+public:
+  S4(int v):a(v) { }
+};
+class S5 {
+  int a;
+  S5():a(0) {}
+  S5(const S5 &s5):a(s5.a) { }
+public:
+  S5(int v):a(v) { }
+};
+
+S3 h;
+#pragma omp threadprivate(h)
+
+typedef struct {
+  int a;
+} S6;
+
+template <typename T>
+T tmain(T argc) {
+  const T da[5] = { 0 };
+  S6 h[10];
+  auto &rh = h;
+  T i;
+  T &j = i;
+  T *k = &j;
+  T *&z = k;
+  T aa[10];
+  auto &raa = aa;
+#pragma omp target has_device_addr(k)
+  {}
+#pragma omp target has_device_addr(z)
+  {}
+#pragma omp target has_device_addr(aa)
+  {}
+#pragma omp target has_device_addr(raa)
+  {}
+#pragma omp target has_device_addr(h)
+  {}
+#pragma omp target has_device_addr(rh)
+  {}
+#pragma omp target has_device_addr(da)
+  {}
+  return 0;
+}
+
+// CHECK: template<> int tmain<int>(int argc) {
+// CHECK-NEXT: const int da[5] = {0};
+// CHECK-NEXT: S6 h[10];
+// CHECK-NEXT: auto &rh = h;
+// CHECK-NEXT: int i;
+// CHECK-NEXT: int &j = i;
+// CHECK-NEXT: int *k = &j;
+// CHECK-NEXT: int *&z = k;
+// CHECK-NEXT: int aa[10];
+// CHECK-NEXT: auto &raa = aa;
+// CHECK-NEXT: #pragma omp target has_device_addr(k)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(z)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(aa)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(raa)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(h)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(rh)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(da)
+
+// CHECK: template<> int *tmain<int *>(int *argc) {
+// CHECK-NEXT: int *const da[5] = {0};
+// CHECK-NEXT: S6 h[10];
+// CHECK-NEXT: auto &rh = h;
+// CHECK-NEXT: int *i;
+// CHECK-NEXT: int *&j = i;
+// CHECK-NEXT: int **k = &j;
+// CHECK-NEXT: int **&z = k;
+// CHECK-NEXT: int *aa[10];
+// CHECK-NEXT: auto &raa = aa;
+// CHECK-NEXT: #pragma omp target has_device_addr(k)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(z)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(aa)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(raa)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(h)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(rh)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(da)
+
+// CHECK-LABEL: int main(int argc, char **argv) {
+int main(int argc, char **argv) {
+  const int da[5] = { 0 };
+  S6 h[10];
+  auto &rh = h;
+  int i;
+  int &j = i;
+  int *k = &j;
+  int *&z = k;
+  int aa[10];
+  auto &raa = aa;
+// CHECK-NEXT: const int da[5] = {0};
+// CHECK-NEXT: S6 h[10];
+// CHECK-NEXT: auto &rh = h;
+// CHECK-NEXT: int i;
+// CHECK-NEXT: int &j = i;
+// CHECK-NEXT: int *k = &j;
+// CHECK-NEXT: int *&z = k;
+// CHECK-NEXT: int aa[10];
+// CHECK-NEXT: auto &raa = aa;
+#pragma omp target has_device_addr(k)
+// CHECK-NEXT: #pragma omp target has_device_addr(k)
+  {}
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+#pragma omp target has_device_addr(z)
+// CHECK-NEXT: #pragma omp target has_device_addr(z)
+  {}
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+#pragma omp target has_device_addr(aa)
+// CHECK-NEXT: #pragma omp target has_device_addr(aa)
+  {}
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+#pragma omp target has_device_addr(raa)
+// CHECK-NEXT: #pragma omp target has_device_addr(raa)
+  {}
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+#pragma omp target has_device_addr(h)
+// CHECK-NEXT: #pragma omp target has_device_addr(h)
+  {}
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+#pragma omp target has_device_addr(rh)
+// CHECK-NEXT: #pragma omp target has_device_addr(rh)
+  {}
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+#pragma omp target has_device_addr(da)
+// CHECK-NEXT: #pragma omp target has_device_addr(da)
+  {}
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target has_device_addr(da[1:3])
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+#pragma omp target has_device_addr(da[1:3])
+  {}
+  return tmain<int>(argc) + *tmain<int *>(&argc);
+}
+
+struct SomeKernel {
+  int targetDev;
+  float devPtr;
+  SomeKernel();
+  ~SomeKernel();
+
+  template<unsigned int nRHS>
+  void apply() {
+    #pragma omp target has_device_addr(devPtr) device(targetDev)
+    {
+    }
+// CHECK:  #pragma omp target has_device_addr(this->devPtr) device(this->targetDev)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+  }
+// CHECK: template<> void apply<32U>() {
+// CHECK: #pragma omp target has_device_addr(this->devPtr) device(this->targetDev)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+};
+
+void use_template() {
+  SomeKernel aKern;
+  aKern.apply<32>();
+}
+#endif
diff --git a/clang/test/OpenMP/target_has_device_addr_messages.cpp b/clang/test/OpenMP/target_has_device_addr_messages.cpp
new file mode 100644
--- /dev/null
+++ b/clang/test/OpenMP/target_has_device_addr_messages.cpp
@@ -0,0 +1,273 @@
+// RUN: %clang_cc1 -std=c++11 -fopenmp-version=51 -verify \
+// RUN:  -fopenmp -ferror-limit 200 %s -Wuninitialized
+
+// RUN: %clang_cc1 -std=c++11 -fopenmp-version=51 -verify \
+// RUN:  -fopenmp-simd -ferror-limit 200 %s -Wuninitialized
+
+struct ST {
+  int *a;
+};
+typedef int arr[10];
+typedef ST STarr[10];
+struct SA {
+  const int d = 5;
+  const int da[5] = { 0 };
+  ST e;
+  ST g[10];
+  STarr &rg = g;
+  int i;
+  int &j = i;
+  int *k = &j;
+  int *&z = k;
+  int aa[10];
+  arr &raa = aa;
+  void func(int arg) {
+#pragma omp target has_device_addr // expected-error {{expected '(' after 'has_device_addr'}}
+    {}
+#pragma omp target has_device_addr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+    {}
+#pragma omp target has_device_addr() // expected-error {{expected expression}}
+    {}
+#pragma omp target has_device_addr(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
+    {}
+#pragma omp target has_device_addr(arg // expected-error {{expected ')'}} expected-note {{to match this '('}}
+    {}
+#pragma omp target has_device_addr(i) // OK
+    {}
+#pragma omp target has_device_addr(j) // OK
+    {}
+#pragma omp target has_device_addr(k) // OK
+    {}
+#pragma omp target has_device_addr(z) // OK
+    {}
+#pragma omp target has_device_addr(aa) // OK
+    {}
+#pragma omp target has_device_addr(raa) // OK
+    {}
+#pragma omp target has_device_addr(e) // OK
+    {}
+#pragma omp target has_device_addr(g) // OK
+    {}
+#pragma omp target has_device_addr(rg) // OK
+    {}
+#pragma omp target has_device_addr(k,i,j) // OK
+    {}
+#pragma omp target has_device_addr(d) // OK
+    {}
+#pragma omp target has_device_addr(da) // OK
+    {}
+  return;
+ }
+};
+struct SB {
+  unsigned A;
+  unsigned B;
+  float Arr[100];
+  float *Ptr;
+  float *foo() {
+    return &Arr[0];
+  }
+};
+
+struct SC {
+  unsigned A : 2;
+  unsigned B : 3;
+  unsigned C;
+  unsigned D;
+  float Arr[100];
+  SB S;
+  SB ArrS[100];
+  SB *PtrS;
+  SB *&RPtrS;
+  float *Ptr;
+
+  SC(SB *&_RPtrS) : RPtrS(_RPtrS) {}
+};
+
+union SD {
+  unsigned A;
+  float B;
+};
+
+struct S1;
+extern S1 a;
+class S2 {
+  mutable int a;
+public:
+  S2():a(0) { }
+  S2(S2 &s2):a(s2.a) { }
+  static float S2s;
+  static const float S2sc;
+};
+const float S2::S2sc = 0;
+const S2 b;
+const S2 ba[5];
+class S3 {
+  int a;
+public:
+  S3():a(0) { }
+  S3(S3 &s3):a(s3.a) { }
+};
+const S3 c;
+const S3 ca[5];
+extern const int f;
+class S4 {
+  int a;
+  S4();
+  S4(const S4 &s4);
+public:
+  S4(int v):a(v) { }
+};
+class S5 {
+  int a;
+  S5():a(0) {}
+  S5(const S5 &s5):a(s5.a) { }
+public:
+  S5(int v):a(v) { }
+};
+
+S3 h;
+#pragma omp threadprivate(h)
+
+typedef struct {
+  int a;
+} S6;
+
+template <typename T, int I>
+T tmain(T argc) {
+  const T d = 5;
+  const T da[5] = { 0 };
+  S4 e(4);
+  S5 g(5);
+  S6 h[10];
+  auto &rh = h;
+  T i;
+  T &j = i;
+  T *k = &j;
+  T *&z = k;
+  T aa[10];
+  auto &raa = aa;
+  S6 *ps;
+#pragma omp target has_device_addr // expected-error {{expected '(' after 'has_device_addr'}}
+  {}
+#pragma omp target has_device_addr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+  {}
+#pragma omp target has_device_addr() // expected-error {{expected expression}}
+  {}
+#pragma omp target has_device_addr(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
+  {}
+#pragma omp target has_device_addr(argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+  {}
+#pragma omp target has_device_addr(i) // OK
+  {}
+#pragma omp target has_device_addr(j) // OK
+  {}
+#pragma omp target has_device_addr(k) // OK
+  {}
+#pragma omp target has_device_addr(z) // OK
+  {}
+#pragma omp target has_device_addr(aa) // OK
+  {}
+#pragma omp target has_device_addr(raa) // OK
+  {}
+#pragma omp target has_device_addr(e) // OK
+  {}
+#pragma omp target has_device_addr(g) // OK
+  {}
+#pragma omp target has_device_addr(h) // OK
+  {}
+#pragma omp target has_device_addr(rh) // OK
+  {}
+#pragma omp target has_device_addr(k,i,j) // OK
+  {}
+#pragma omp target has_device_addr(d) // OK
+  {}
+#pragma omp target has_device_addr(da) // OK
+  {}
+#pragma omp target map(ps) has_device_addr(ps) // expected-error{{variable already marked as mapped in current construct}} expected-note{{used here}}
+  {}
+#pragma omp target has_device_addr(ps) map(ps) // expected-error{{variable already marked as mapped in current construct}} expected-note{{used here}}
+  {}
+#pragma omp target map(ps->a) has_device_addr(ps) // expected-error{{variable already marked as mapped in current construct}} expected-note{{used here}}
+  {}
+#pragma omp target has_device_addr(ps) map(ps->a) // expected-error{{pointer cannot be mapped along with a section derived from itself}} expected-note{{used here}}
+  {}
+#pragma omp target has_device_addr(ps) firstprivate(ps) // expected-error{{firstprivate variable cannot be in a has_device_addr clause in '#pragma omp target' directive}}
+  {}
+#pragma omp target firstprivate(ps) has_device_addr(ps) // expected-error{{firstprivate variable cannot be in a has_device_addr clause in '#pragma omp target' directive}} expected-note{{defined as firstprivate}}
+  {}
+#pragma omp target has_device_addr(ps) private(ps) // expected-error{{private variable cannot be in a has_device_addr clause in '#pragma omp target' directive}}
+  {}
+#pragma omp target private(ps) has_device_addr(ps) // expected-error{{private variable cannot be in a has_device_addr clause in '#pragma omp target' directive}} expected-note{{defined as private}}
+  {}
+  return 0;
+}
+
+int main(int argc, char **argv) {
+  const int d = 5;
+  const int da[5] = { 0 };
+  S4 e(4);
+  S5 g(5);
+  S6 h[10];
+  auto &rh = h;
+  int i;
+  int &j = i;
+  int *k = &j;
+  int *&z = k;
+  int aa[10];
+  auto &raa = aa;
+  S6 *ps;
+#pragma omp target has_device_addr // expected-error {{expected '(' after 'has_device_addr'}}
+  {}
+#pragma omp target has_device_addr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+  {}
+#pragma omp target has_device_addr() // expected-error {{expected expression}}
+  {}
+#pragma omp target has_device_addr(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
+  {}
+#pragma omp target has_device_addr(argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+  {}
+#pragma omp target has_device_addr(i) // OK
+  {}
+#pragma omp target has_device_addr(j) // OK
+  {}
+#pragma omp target has_device_addr(k) // OK
+  {}
+#pragma omp target has_device_addr(z) // OK
+  {}
+#pragma omp target has_device_addr(aa) // OK
+  {}
+#pragma omp target has_device_addr(raa) // OK
+  {}
+#pragma omp target has_device_addr(e) // OK
+  {}
+#pragma omp target has_device_addr(g) // OK
+  {}
+#pragma omp target has_device_addr(h) // OK
+  {}
+#pragma omp target has_device_addr(rh) // OK
+  {}
+#pragma omp target has_device_addr(k,i,j) // OK
+  {}
+#pragma omp target has_device_addr(d) // OK
+  {}
+#pragma omp target has_device_addr(da) // OK
+  {}
+#pragma omp target map(ps) has_device_addr(ps) // expected-error{{variable already marked as mapped in current construct}} expected-note{{used here}}
+  {}
+#pragma omp target has_device_addr(ps) map(ps) // expected-error{{variable already marked as mapped in current construct}} expected-note{{used here}}
+  {}
+#pragma omp target map(ps->a) has_device_addr(ps) // expected-error{{variable already marked as mapped in current construct}} expected-note{{used here}}
+  {}
+#pragma omp target has_device_addr(ps) map(ps->a) // expected-error{{pointer cannot be mapped along with a section derived from itself}} expected-note{{used here}}
+  {}
+#pragma omp target has_device_addr(ps) firstprivate(ps) // expected-error{{firstprivate variable cannot be in a has_device_addr clause in '#pragma omp target' directive}}
+  {}
+#pragma omp target firstprivate(ps) has_device_addr(ps) // expected-error{{firstprivate variable cannot be in a has_device_addr clause in '#pragma omp target' directive}} expected-note{{defined as firstprivate}}
+  {}
+#pragma omp target has_device_addr(ps) private(ps) // expected-error{{private variable cannot be in a has_device_addr clause in '#pragma omp target' directive}}
+  {}
+#pragma omp target private(ps) has_device_addr(ps) // expected-error{{private variable cannot be in a has_device_addr clause in '#pragma omp target' directive}} expected-note{{defined as private}}
+  {}
+  return tmain<int, 3>(argc);
+}
diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -2573,6 +2573,10 @@
     const OMPIsDevicePtrClause *C) {
   VisitOMPClauseList(C);
 }
+void OMPClauseEnqueue::VisitOMPHasDeviceAddrClause(
+    const OMPHasDeviceAddrClause *C) {
+  VisitOMPClauseList(C);
+}
 void OMPClauseEnqueue::VisitOMPNontemporalClause(
     const OMPNontemporalClause *C) {
   VisitOMPClauseList(C);
diff --git a/flang/lib/Semantics/check-omp-structure.cpp b/flang/lib/Semantics/check-omp-structure.cpp
--- a/flang/lib/Semantics/check-omp-structure.cpp
+++ b/flang/lib/Semantics/check-omp-structure.cpp
@@ -1683,6 +1683,7 @@
 CHECK_SIMPLE_CLAUSE(Threads, OMPC_threads)
 CHECK_SIMPLE_CLAUSE(Inbranch, OMPC_inbranch)
 CHECK_SIMPLE_CLAUSE(IsDevicePtr, OMPC_is_device_ptr)
+CHECK_SIMPLE_CLAUSE(HasDeviceAddr, OMPC_has_device_addr)
 CHECK_SIMPLE_CLAUSE(Link, OMPC_link)
 CHECK_SIMPLE_CLAUSE(Indirect, OMPC_indirect)
 CHECK_SIMPLE_CLAUSE(Mergeable, OMPC_mergeable)
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMP.td b/llvm/include/llvm/Frontend/OpenMP/OMP.td
--- a/llvm/include/llvm/Frontend/OpenMP/OMP.td
+++ b/llvm/include/llvm/Frontend/OpenMP/OMP.td
@@ -253,6 +253,11 @@
   let flangClass = "Name";
   let isValueList = true;
 }
+def OMPC_HasDeviceAddr : Clause<"has_device_addr"> {
+  let clangClass = "OMPHasDeviceAddrClause";
+  let flangClass = "Name";
+  let isValueList = true;
+}
 def OMPC_TaskReduction : Clause<"task_reduction"> {
   let clangClass = "OMPTaskReductionClause";
   let flangClass = "OmpReductionClause";
@@ -556,6 +561,7 @@
     VersionedClause<OMPC_Depend>,
     VersionedClause<OMPC_FirstPrivate>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_Reduction>,
     VersionedClause<OMPC_Allocate>,
     VersionedClause<OMPC_UsesAllocators, 50>
@@ -652,6 +658,7 @@
     VersionedClause<OMPC_Shared>,
     VersionedClause<OMPC_Reduction>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_Allocate>,
     VersionedClause<OMPC_UsesAllocators, 50>
   ];
@@ -684,6 +691,7 @@
     VersionedClause<OMPC_Ordered>,
     VersionedClause<OMPC_Linear>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_Allocate>,
     VersionedClause<OMPC_Order, 50>,
     VersionedClause<OMPC_UsesAllocators, 50>
@@ -700,6 +708,7 @@
     VersionedClause<OMPC_Reduction>,
     VersionedClause<OMPC_Linear>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_Allocator>,
     VersionedClause<OMPC_Order>,
     VersionedClause<OMPC_UsesAllocators>,
@@ -1133,6 +1142,7 @@
     VersionedClause<OMPC_SimdLen>,
     VersionedClause<OMPC_Aligned>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_Allocate>,
     VersionedClause<OMPC_NonTemporal, 50>,
     VersionedClause<OMPC_Order, 50>,
@@ -1163,6 +1173,7 @@
     VersionedClause<OMPC_SimdLen>,
     VersionedClause<OMPC_Aligned>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_Allocate>,
     VersionedClause<OMPC_NonTemporal>,
     VersionedClause<OMPC_Order>,
@@ -1176,6 +1187,7 @@
     VersionedClause<OMPC_Depend>,
     VersionedClause<OMPC_FirstPrivate>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_LastPrivate>,
     VersionedClause<OMPC_Linear>,
     VersionedClause<OMPC_Map>,
@@ -1349,6 +1361,7 @@
     VersionedClause<OMPC_Depend>,
     VersionedClause<OMPC_FirstPrivate>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_Reduction>,
     VersionedClause<OMPC_Allocate>,
     VersionedClause<OMPC_UsesAllocators, 50>,
@@ -1372,6 +1385,7 @@
     VersionedClause<OMPC_Depend>,
     VersionedClause<OMPC_FirstPrivate>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_Reduction>,
     VersionedClause<OMPC_Allocate>,
     VersionedClause<OMPC_UsesAllocators, 50>,
@@ -1402,6 +1416,7 @@
     VersionedClause<OMPC_DefaultMap>,
     VersionedClause<OMPC_FirstPrivate>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_Default>,
     VersionedClause<OMPC_Shared>,
     VersionedClause<OMPC_Reduction>,
@@ -1427,6 +1442,7 @@
     VersionedClause<OMPC_Depend>,
     VersionedClause<OMPC_FirstPrivate>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_Reduction>,
     VersionedClause<OMPC_Allocate>,
     VersionedClause<OMPC_UsesAllocators>,
@@ -1463,6 +1479,7 @@
     VersionedClause<OMPC_DefaultMap>,
     VersionedClause<OMPC_FirstPrivate>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_Default>,
     VersionedClause<OMPC_Shared>,
     VersionedClause<OMPC_Reduction>,
@@ -1492,6 +1509,7 @@
     VersionedClause<OMPC_Depend>,
     VersionedClause<OMPC_FirstPrivate>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_Reduction>,
     VersionedClause<OMPC_Allocate>,
     VersionedClause<OMPC_UsesAllocators>,
@@ -1530,6 +1548,7 @@
     VersionedClause<OMPC_FirstPrivate>,
     VersionedClause<OMPC_If>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_LastPrivate>,
     VersionedClause<OMPC_Linear>,
     VersionedClause<OMPC_Map>,
@@ -1741,6 +1760,7 @@
   let allowedClauses = [
     VersionedClause<OMPC_Device>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_NoWait>,
     VersionedClause<OMPC_Depend>,
     VersionedClause<OMPC_Novariants>,
@@ -1790,6 +1810,7 @@
     VersionedClause<OMPC_Device>,
     VersionedClause<OMPC_FirstPrivate>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_LastPrivate>,
     VersionedClause<OMPC_Map>,
     VersionedClause<OMPC_Private>,
@@ -1836,6 +1857,7 @@
     VersionedClause<OMPC_Device>,
     VersionedClause<OMPC_FirstPrivate>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_LastPrivate>,
     VersionedClause<OMPC_Map>,
     VersionedClause<OMPC_Private>,