[OPENMP50]Initial support for use_device_addr clause.

[ABataev]

Summary

Added parsing/sema analysis/serialization support for use_device_addr
clauses.

[ABataev]

Rebase + cleanup

[jdoerfert]

LGTM, one nit below

clang/lib/Sema/SemaOpenMP.cpp

10179

The restriction text quoted above doesn't seem to match the code anymore (also before).

clang/include/clang/AST/OpenMPClause.h

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_use_device_ptr;
   }
 };
 
+/// This represents clause 'use_device_addr' in the '#pragma omp ...'
+/// directives.
+///
+/// \code
+/// #pragma omp target data use_device_addr(a,b)
+/// \endcode
+/// In this example directive '#pragma omp target data' has clause
+/// 'use_device_addr' with the variables 'a' and 'b'.
+class OMPUseDeviceAddrClause final
+    : public OMPMappableExprListClause<OMPUseDeviceAddrClause>,
+      private llvm::TrailingObjects<
+          OMPUseDeviceAddrClause, 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 OMPUseDeviceAddrClause(const OMPVarListLocTy &Locs,
+                                  const OMPMappableExprListSizeTy &Sizes)
+      : OMPMappableExprListClause(llvm::omp::OMPC_use_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 OMPUseDeviceAddrClause(const OMPMappableExprListSizeTy &Sizes)
+      : OMPMappableExprListClause(llvm::omp::OMPC_use_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 3 * varlist_size();
+  }
+  size_t numTrailingObjects(OverloadToken<ValueDecl *>) const {
+    return getUniqueDeclarationsNum();
+  }
+  size_t numTrailingObjects(OverloadToken<unsigned>) const {
+    return getUniqueDeclarationsNum() + getTotalComponentListNum();
+  }
+
+  /// Sets the list of references to private copies with initializers for new
+  /// private variables.
+  /// \param VL List of references.
+  void setPrivateCopies(ArrayRef<Expr *> VL);
+
+  /// Gets the list of references to private copies with initializers for new
+  /// private variables.
+  MutableArrayRef<Expr *> getPrivateCopies() {
+    return MutableArrayRef<Expr *>(varlist_end(), varlist_size());
+  }
+  ArrayRef<const Expr *> getPrivateCopies() const {
+    return llvm::makeArrayRef(varlist_end(), varlist_size());
+  }
+
+  /// Sets the list of references to initializer variables for new private
+  /// variables.
+  /// \param VL List of references.
+  void setInits(ArrayRef<Expr *> VL);
+
+  /// Gets the list of references to initializer variables for new private
+  /// variables.
+  MutableArrayRef<Expr *> getInits() {
+    return MutableArrayRef<Expr *>(getPrivateCopies().end(), varlist_size());
+  }
+  ArrayRef<const Expr *> getInits() const {
+    return llvm::makeArrayRef(getPrivateCopies().end(), varlist_size());
+  }
+
+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 PrivateVars Expressions referring to private copies.
+  /// \param Inits Expressions referring to private copy initializers.
+  /// \param Declarations Declarations used in the clause.
+  /// \param ComponentLists Component lists used in the clause.
+  static OMPUseDeviceAddrClause *
+  Create(const ASTContext &C, const OMPVarListLocTy &Locs,
+         ArrayRef<Expr *> Vars, ArrayRef<Expr *> PrivateVars,
+         ArrayRef<Expr *> Inits, 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 OMPUseDeviceAddrClause *
+  CreateEmpty(const ASTContext &C, const OMPMappableExprListSizeTy &Sizes);
+
+  using private_copies_iterator = MutableArrayRef<Expr *>::iterator;
+  using private_copies_const_iterator = ArrayRef<const Expr *>::iterator;
+  using private_copies_range = llvm::iterator_range<private_copies_iterator>;
+  using private_copies_const_range =
+      llvm::iterator_range<private_copies_const_iterator>;
+
+  private_copies_range private_copies() {
+    return private_copies_range(getPrivateCopies().begin(),
+                                getPrivateCopies().end());
+  }
+
+  private_copies_const_range private_copies() const {
+    return private_copies_const_range(getPrivateCopies().begin(),
+                                      getPrivateCopies().end());
+  }
+
+  using inits_iterator = MutableArrayRef<Expr *>::iterator;
+  using inits_const_iterator = ArrayRef<const Expr *>::iterator;
+  using inits_range = llvm::iterator_range<inits_iterator>;
+  using inits_const_range = llvm::iterator_range<inits_const_iterator>;
+
+  inits_range inits() {
+    return inits_range(getInits().begin(), getInits().end());
+  }
+
+  inits_const_range inits() const {
+    return inits_const_range(getInits().begin(), getInits().end());
+  }
+
+  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<OMPUseDeviceAddrClause *>(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_use_device_addr;
+  }
+};
+
 /// This represents clause 'is_device_ptr' in the '#pragma omp ...'
 /// directives.
 ///
 /// \code
 /// #pragma omp target is_device_ptr(a,b)
 /// \endcode
 /// In this example directive '#pragma omp target' has clause
 /// 'is_device_ptr' with the variables 'a' and 'b'.

clang/include/clang/AST/RecursiveASTVisitor.h

 template <typename Derived>
 bool RecursiveASTVisitor<Derived>::VisitOMPUseDevicePtrClause(
     OMPUseDevicePtrClause *C) {
   TRY_TO(VisitOMPClauseList(C));
   return true;
 }
 
 template <typename Derived>
+bool RecursiveASTVisitor<Derived>::VisitOMPUseDeviceAddrClause(
+    OMPUseDeviceAddrClause *C) {
+  TRY_TO(VisitOMPClauseList(C));
+  return true;
+}
+
+template <typename Derived>
 bool RecursiveASTVisitor<Derived>::VisitOMPIsDevicePtrClause(
     OMPIsDevicePtrClause *C) {
   TRY_TO(VisitOMPClauseList(C));
   return true;
 }
 
 template <typename Derived>
 bool RecursiveASTVisitor<Derived>::VisitOMPNontemporalClause(

clang/include/clang/Sema/Sema.h

public:
   /// Called on well-formed 'from' clause.
   OMPClause *ActOnOpenMPFromClause(
       ArrayRef<Expr *> VarList, CXXScopeSpec &MapperIdScopeSpec,
       DeclarationNameInfo &MapperId, const OMPVarListLocTy &Locs,
       ArrayRef<Expr *> UnresolvedMappers = llvm::None);
   /// Called on well-formed 'use_device_ptr' clause.
   OMPClause *ActOnOpenMPUseDevicePtrClause(ArrayRef<Expr *> VarList,
                                            const OMPVarListLocTy &Locs);
+  /// Called on well-formed 'use_device_addr' clause.
+  OMPClause *ActOnOpenMPUseDeviceAddrClause(ArrayRef<Expr *> VarList,
+                                            const OMPVarListLocTy &Locs);
   /// Called on well-formed 'is_device_ptr' clause.
   OMPClause *ActOnOpenMPIsDevicePtrClause(ArrayRef<Expr *> VarList,
                                           const OMPVarListLocTy &Locs);
   /// Called on well-formed 'nontemporal' clause.
   OMPClause *ActOnOpenMPNontemporalClause(ArrayRef<Expr *> VarList,
                                           SourceLocation StartLoc,
                                           SourceLocation LParenLoc,
                                           SourceLocation EndLoc);

clang/lib/AST/OpenMPClause.cpp

const OMPClauseWithPreInit *OMPClauseWithPreInit::get(const OMPClause *C) {
   case OMPC_nogroup:
   case OMPC_hint:
   case OMPC_defaultmap:
   case OMPC_unknown:
   case OMPC_uniform:
   case OMPC_to:
   case OMPC_from:
   case OMPC_use_device_ptr:
+  case OMPC_use_device_addr:
   case OMPC_is_device_ptr:
   case OMPC_unified_address:
   case OMPC_unified_shared_memory:
   case OMPC_reverse_offload:
   case OMPC_dynamic_allocators:
   case OMPC_atomic_default_mem_order:
   case OMPC_device_type:
   case OMPC_match:
const OMPClauseWithPostUpdate *OMPClauseWithPostUpdate::get(const OMPClause *C) {
   case OMPC_num_tasks:
   case OMPC_hint:
   case OMPC_defaultmap:
   case OMPC_unknown:
   case OMPC_uniform:
   case OMPC_to:
   case OMPC_from:
   case OMPC_use_device_ptr:
+  case OMPC_use_device_addr:
   case OMPC_is_device_ptr:
   case OMPC_unified_address:
   case OMPC_unified_shared_memory:
   case OMPC_reverse_offload:
   case OMPC_dynamic_allocators:
   case OMPC_atomic_default_mem_order:
   case OMPC_device_type:
   case OMPC_match:
void *Mem = C.Allocate(
       totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
                        OMPClauseMappableExprCommon::MappableComponent>(
           3 * Sizes.NumVars, Sizes.NumUniqueDeclarations,
           Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
           Sizes.NumComponents));
   return new (Mem) OMPUseDevicePtrClause(Sizes);
 }
 
+void OMPUseDeviceAddrClause::setPrivateCopies(ArrayRef<Expr *> VL) {
+  assert(VL.size() == varlist_size() &&
+         "Number of private copies is not the same as the preallocated buffer");
+  std::copy(VL.begin(), VL.end(), varlist_end());
+}
+
+void OMPUseDeviceAddrClause::setInits(ArrayRef<Expr *> VL) {
+  assert(VL.size() == varlist_size() &&
+         "Number of inits is not the same as the preallocated buffer");
+  std::copy(VL.begin(), VL.end(), getPrivateCopies().end());
+}
+
+OMPUseDeviceAddrClause *OMPUseDeviceAddrClause::Create(
+    const ASTContext &C, const OMPVarListLocTy &Locs, ArrayRef<Expr *> Vars,
+    ArrayRef<Expr *> PrivateVars, ArrayRef<Expr *> Inits,
+    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:
+  // 3 x NumVars x Expr* - we have an original list expression for each clause
+  // list entry and an equal number of private copies and inits.
+  // 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>(
+          3 * Sizes.NumVars, Sizes.NumUniqueDeclarations,
+          Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
+          Sizes.NumComponents));
+
+  auto *Clause = new (Mem) OMPUseDeviceAddrClause(Locs, Sizes);
+
+  Clause->setVarRefs(Vars);
+  Clause->setPrivateCopies(PrivateVars);
+  Clause->setInits(Inits);
+  Clause->setClauseInfo(Declarations, ComponentLists);
+  return Clause;
+}
+
+OMPUseDeviceAddrClause *
+OMPUseDeviceAddrClause::CreateEmpty(const ASTContext &C,
+                                    const OMPMappableExprListSizeTy &Sizes) {
+  void *Mem = C.Allocate(
+      totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
+                       OMPClauseMappableExprCommon::MappableComponent>(
+          3 * Sizes.NumVars, Sizes.NumUniqueDeclarations,
+          Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
+          Sizes.NumComponents));
+  return new (Mem) OMPUseDeviceAddrClause(Sizes);
+}
+
 OMPIsDevicePtrClause *
 OMPIsDevicePtrClause::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);

 void OMPClausePrinter::VisitOMPUseDevicePtrClause(OMPUseDevicePtrClause *Node) {
   if (!Node->varlist_empty()) {
     OS << "use_device_ptr";
     VisitOMPClauseList(Node, '(');
     OS << ")";
   }
 }
 
+void OMPClausePrinter::VisitOMPUseDeviceAddrClause(
+    OMPUseDeviceAddrClause *Node) {
+  if (!Node->varlist_empty()) {
+    OS << "use_device_addr";
+    VisitOMPClauseList(Node, '(');
+    OS << ")";
+  }
+}
+
 void OMPClausePrinter::VisitOMPIsDevicePtrClause(OMPIsDevicePtrClause *Node) {
   if (!Node->varlist_empty()) {
     OS << "is_device_ptr";
     VisitOMPClauseList(Node, '(');
     OS << ")";
   }
 }
 

clang/lib/AST/StmtProfile.cpp

 }
 void OMPClauseProfiler::VisitOMPFromClause(const OMPFromClause *C) {
   VisitOMPClauseList(C);
 }
 void OMPClauseProfiler::VisitOMPUseDevicePtrClause(
     const OMPUseDevicePtrClause *C) {
   VisitOMPClauseList(C);
 }
+void OMPClauseProfiler::VisitOMPUseDeviceAddrClause(
+    const OMPUseDeviceAddrClause *C) {
+  VisitOMPClauseList(C);
+}
 void OMPClauseProfiler::VisitOMPIsDevicePtrClause(
     const OMPIsDevicePtrClause *C) {
   VisitOMPClauseList(C);
 }
 void OMPClauseProfiler::VisitOMPNontemporalClause(
     const OMPNontemporalClause *C) {
   VisitOMPClauseList(C);
   for (auto *E : C->private_refs())

clang/lib/Basic/OpenMPKinds.cpp

#include "clang/Basic/OpenMPKinds.def"
   case OMPC_thread_limit:
   case OMPC_priority:
   case OMPC_grainsize:
   case OMPC_nogroup:
   case OMPC_num_tasks:
   case OMPC_hint:
   case OMPC_uniform:
   case OMPC_use_device_ptr:
+  case OMPC_use_device_addr:
   case OMPC_is_device_ptr:
   case OMPC_unified_address:
   case OMPC_unified_shared_memory:
   case OMPC_reverse_offload:
   case OMPC_dynamic_allocators:
   case OMPC_match:
   case OMPC_nontemporal:
   case OMPC_destroy:
#include "clang/Basic/OpenMPKinds.def"
   case OMPC_thread_limit:
   case OMPC_priority:
   case OMPC_grainsize:
   case OMPC_nogroup:
   case OMPC_num_tasks:
   case OMPC_hint:
   case OMPC_uniform:
   case OMPC_use_device_ptr:
+  case OMPC_use_device_addr:
   case OMPC_is_device_ptr:
   case OMPC_unified_address:
   case OMPC_unified_shared_memory:
   case OMPC_reverse_offload:
   case OMPC_dynamic_allocators:
   case OMPC_match:
   case OMPC_nontemporal:
   case OMPC_destroy:

clang/lib/CodeGen/CGStmtOpenMP.cpp

static void emitOMPAtomicExpr(CodeGenFunction &CGF, OpenMPClauseKind Kind,
   case OMPC_num_tasks:
   case OMPC_hint:
   case OMPC_dist_schedule:
   case OMPC_defaultmap:
   case OMPC_uniform:
   case OMPC_to:
   case OMPC_from:
   case OMPC_use_device_ptr:
+  case OMPC_use_device_addr:
   case OMPC_is_device_ptr:
   case OMPC_unified_address:
   case OMPC_unified_shared_memory:
   case OMPC_reverse_offload:
   case OMPC_dynamic_allocators:
   case OMPC_atomic_default_mem_order:
   case OMPC_device_type:
   case OMPC_match:

clang/lib/Parse/ParseOpenMP.cpp

 ///       update-clause | capture-clause | seq_cst-clause | device-clause |
 ///       simdlen-clause | threads-clause | simd-clause | num_teams-clause |
 ///       thread_limit-clause | priority-clause | grainsize-clause |
 ///       nogroup-clause | num_tasks-clause | hint-clause | to-clause |
 ///       from-clause | is_device_ptr-clause | task_reduction-clause |
 ///       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
+///       exclusive-clause | uses_allocators-clause | use_device_addr-clause
 ///
 OMPClause *Parser::ParseOpenMPClause(OpenMPDirectiveKind DKind,
                                      OpenMPClauseKind CKind, bool FirstClause) {
   OMPClause *Clause = nullptr;
   bool ErrorFound = false;
   bool WrongDirective = false;
   // Check if clause is allowed for the given directive.
   if (CKind != OMPC_unknown &&
OMPClause *Parser::ParseOpenMPClause(OpenMPDirectiveKind DKind,
   case OMPC_copyin:
   case OMPC_copyprivate:
   case OMPC_flush:
   case OMPC_depend:
   case OMPC_map:
   case OMPC_to:
   case OMPC_from:
   case OMPC_use_device_ptr:
+  case OMPC_use_device_addr:
   case OMPC_is_device_ptr:
   case OMPC_allocate:
   case OMPC_nontemporal:
   case OMPC_inclusive:
   case OMPC_exclusive:
   case OMPC_affinity:
     Clause = ParseOpenMPVarListClause(DKind, CKind, WrongDirective);
     break;

 ///          [ mapper '(' mapper-identifier ')' [,] ]
 ///          to | from | tofrom | alloc | release | delete ':' ] list ')';
 ///    to-clause:
 ///       'to' '(' [ mapper '(' mapper-identifier ')' ':' ] list ')'
 ///    from-clause:
 ///       'from' '(' [ mapper '(' mapper-identifier ')' ':' ] list ')'
 ///    use_device_ptr-clause:
 ///       'use_device_ptr' '(' list ')'
+///    use_device_addr-clause:
+///       'use_device_addr' '(' list ')'
 ///    is_device_ptr-clause:
 ///       'is_device_ptr' '(' list ')'
 ///    allocate-clause:
 ///       'allocate' '(' [ allocator ':' ] list ')'
 ///    nontemporal-clause:
 ///       'nontemporal' '(' list ')'
 ///    inclusive-clause:
 ///       'inclusive' '(' list ')'

clang/lib/Sema/SemaOpenMP.cpp

for (OMPClause *C : Clauses) {
       case OMPC_threads:
       case OMPC_simd:
       case OMPC_map:
       case OMPC_nogroup:
       case OMPC_defaultmap:
       case OMPC_to:
       case OMPC_from:
       case OMPC_use_device_ptr:
+      case OMPC_use_device_addr:
       case OMPC_is_device_ptr:
       case OMPC_nontemporal:
       case OMPC_order:
       case OMPC_destroy:
       case OMPC_inclusive:
       case OMPC_exclusive:
       case OMPC_uses_allocators:
       case OMPC_affinity:
StmtResult Sema::ActOnOpenMPTargetDataDirective(ArrayRef<OMPClause *> Clauses,
                                                 SourceLocation EndLoc) {
   if (!AStmt)
     return StmtError();
 
   assert(isa<CapturedStmt>(AStmt) && "Captured statement expected");
 
   // OpenMP [2.10.1, Restrictions, p. 97]
   // At least one map clause must appear on the directive.
-  if (!hasClauses(Clauses, OMPC_map, OMPC_use_device_ptr)) {
+  if (!hasClauses(Clauses, OMPC_map, OMPC_use_device_ptr) &&
+      (LangOpts.OpenMP < 50 || !hasClauses(Clauses, OMPC_use_device_addr))) {
+    StringRef Expected;
+    if (LangOpts.OpenMP < 50)
+      Expected = "'map' or 'use_device_ptr'";
+    else
+      Expected = "'map', 'use_device_ptr', or 'use_device_addr'";
     Diag(StartLoc, diag::err_omp_no_clause_for_directive)
-        << "'map' or 'use_device_ptr'"
-        << getOpenMPDirectiveName(OMPD_target_data);
+        << Expected << getOpenMPDirectiveName(OMPD_target_data);
     return StmtError();
   }
 
   setFunctionHasBranchProtectedScope();
 
   return OMPTargetDataDirective::Create(Context, StartLoc, EndLoc, Clauses,
                                         AStmt);
 }
OMPClause *Sema::ActOnOpenMPSingleExprClause(OpenMPClauseKind Kind, Expr *Expr,
   case OMPC_nogroup:
   case OMPC_dist_schedule:
   case OMPC_defaultmap:
   case OMPC_unknown:
   case OMPC_uniform:
   case OMPC_to:
   case OMPC_from:
   case OMPC_use_device_ptr:
+  case OMPC_use_device_addr:
   case OMPC_is_device_ptr:
   case OMPC_unified_address:
   case OMPC_unified_shared_memory:
   case OMPC_reverse_offload:
   case OMPC_dynamic_allocators:
   case OMPC_atomic_default_mem_order:
   case OMPC_device_type:
   case OMPC_match:
static OpenMPDirectiveKind getOpenMPCaptureRegionForClause(
   case OMPC_nogroup:
   case OMPC_hint:
   case OMPC_defaultmap:
   case OMPC_unknown:
   case OMPC_uniform:
   case OMPC_to:
   case OMPC_from:
   case OMPC_use_device_ptr:
+  case OMPC_use_device_addr:
   case OMPC_is_device_ptr:
   case OMPC_unified_address:
   case OMPC_unified_shared_memory:
   case OMPC_reverse_offload:
   case OMPC_dynamic_allocators:
   case OMPC_atomic_default_mem_order:
   case OMPC_device_type:
   case OMPC_match:
OMPClause *Sema::ActOnOpenMPSimpleClause(
   case OMPC_hint:
   case OMPC_dist_schedule:
   case OMPC_defaultmap:
   case OMPC_unknown:
   case OMPC_uniform:
   case OMPC_to:
   case OMPC_from:
   case OMPC_use_device_ptr:
+  case OMPC_use_device_addr:
   case OMPC_is_device_ptr:
   case OMPC_unified_address:
   case OMPC_unified_shared_memory:
   case OMPC_reverse_offload:
   case OMPC_dynamic_allocators:
   case OMPC_device_type:
   case OMPC_match:
   case OMPC_nontemporal:
OMPClause *Sema::ActOnOpenMPSingleExprWithArgClause(
   case OMPC_nogroup:
   case OMPC_num_tasks:
   case OMPC_hint:
   case OMPC_unknown:
   case OMPC_uniform:
   case OMPC_to:
   case OMPC_from:
   case OMPC_use_device_ptr:
+  case OMPC_use_device_addr:
   case OMPC_is_device_ptr:
   case OMPC_unified_address:
   case OMPC_unified_shared_memory:
   case OMPC_reverse_offload:
   case OMPC_dynamic_allocators:
   case OMPC_atomic_default_mem_order:
   case OMPC_device_type:
   case OMPC_match:
OMPClause *Sema::ActOnOpenMPClause(OpenMPClauseKind Kind,
   case OMPC_hint:
   case OMPC_dist_schedule:
   case OMPC_defaultmap:
   case OMPC_unknown:
   case OMPC_uniform:
   case OMPC_to:
   case OMPC_from:
   case OMPC_use_device_ptr:
+  case OMPC_use_device_addr:
   case OMPC_is_device_ptr:
   case OMPC_atomic_default_mem_order:
   case OMPC_device_type:
   case OMPC_match:
   case OMPC_nontemporal:
   case OMPC_order:
   case OMPC_detach:
   case OMPC_inclusive:
case OMPC_to:
     break;
   case OMPC_from:
     Res = ActOnOpenMPFromClause(VarList, ReductionOrMapperIdScopeSpec,
                                 ReductionOrMapperId, Locs);
     break;
   case OMPC_use_device_ptr:
     Res = ActOnOpenMPUseDevicePtrClause(VarList, Locs);
     break;
+  case OMPC_use_device_addr:
+    Res = ActOnOpenMPUseDeviceAddrClause(VarList, Locs);
+    break;
   case OMPC_is_device_ptr:
     Res = ActOnOpenMPIsDevicePtrClause(VarList, Locs);
     break;
   case OMPC_allocate:
     Res = ActOnOpenMPAllocateClause(DepModOrTailExpr, VarList, StartLoc,
                                     LParenLoc, ColonLoc, EndLoc);
     break;
   case OMPC_nontemporal:
OMPClause *Sema::ActOnOpenMPUseDevicePtrClause(ArrayRef<Expr *> VarList,
   if (MVLI.ProcessedVarList.empty())
     return nullptr;
 
   return OMPUseDevicePtrClause::Create(
       Context, Locs, MVLI.ProcessedVarList, PrivateCopies, Inits,
       MVLI.VarBaseDeclarations, MVLI.VarComponents);
 }
 
+OMPClause *Sema::ActOnOpenMPUseDeviceAddrClause(ArrayRef<Expr *> VarList,
+                                                const OMPVarListLocTy &Locs) {
+  MappableVarListInfo MVLI(VarList);
+  SmallVector<Expr *, 8> PrivateCopies;
+  SmallVector<Expr *, 8> Inits;
+
+  for (Expr *RefExpr : VarList) {
+    assert(RefExpr && "NULL expr in OpenMP use_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);
+      PrivateCopies.push_back(nullptr);
+      Inits.push_back(nullptr);
+    }
+    ValueDecl *D = Res.first;
+    if (!D)
+      continue;
+
+    QualType Type;
+    auto *OASE = dyn_cast<OMPArraySectionExpr>(SimpleRefExpr);
+    if (OASE) {
+      QualType BaseType =
+          OMPArraySectionExpr::getBaseOriginalType(OASE->getBase());
+      if (const auto *ATy = BaseType->getAsArrayTypeUnsafe())
+        Type = ATy->getElementType();
+      else
+        Type = BaseType->getPointeeType();
+      Type = Type.getNonReferenceType();
+    } else {
+      Type = D->getType().getNonReferenceType();
+    }
+    Type = Context.getPointerType(Type);
+    auto *VD = dyn_cast<VarDecl>(D);
+    // Build the private variable and the expression that refers to it.
+    VarDecl *VDPrivate =
+        buildVarDecl(*this, ELoc, Type, D->getName(),
+                     D->hasAttrs() ? &D->getAttrs() : nullptr,
+                     VD ? cast<DeclRefExpr>(SimpleRefExpr) : nullptr);
+    if (VDPrivate->isInvalidDecl())
+      continue;
+
+    CurContext->addDecl(VDPrivate);
+    DeclRefExpr *VDPrivateRefExpr = buildDeclRefExpr(
+        *this, VDPrivate, RefExpr->getType().getUnqualifiedType(), ELoc);
+
+    // Add temporary variable to initialize the private copy of the pointer.
+    VarDecl *VDInit =
+        buildVarDecl(*this, RefExpr->getExprLoc(), Type, ".devptr.temp");
+    DeclRefExpr *VDInitRefExpr =
+        buildDeclRefExpr(*this, VDInit, Type, RefExpr->getExprLoc());
+    AddInitializerToDecl(VDPrivate,
+                         DefaultLvalueConversion(VDInitRefExpr).get(),
+                         /*DirectInit=*/false);
+
+    // If required, build a capture to implement the privatization initialized
+    // with the current list item value.
+    DeclRefExpr *Ref = nullptr;
+    if (!VD)
+      Ref = buildCapture(*this, D, SimpleRefExpr, /*WithInit=*/true);
+    MVLI.ProcessedVarList.push_back(VD ? RefExpr->IgnoreParens() : Ref);
+    PrivateCopies.push_back(VDPrivateRefExpr);
+    Inits.push_back(VDInitRefExpr);
+
+    // We need to add a data sharing attribute for this variable to make sure it
+    // is correctly captured. A variable that shows up in a use_device_addr has
+    // similar properties of a first private variable.
+    DSAStack->addDSA(D, RefExpr->IgnoreParens(), OMPC_firstprivate, Ref);
+
+    // Create a mappable component for the list item. List items in this clause
+    // only need a component.
+    MVLI.VarBaseDeclarations.push_back(D);
+    MVLI.VarComponents.resize(MVLI.VarComponents.size() + 1);
+    MVLI.VarComponents.back().push_back(
+        OMPClauseMappableExprCommon::MappableComponent(SimpleRefExpr, D));
+  }
+
+  if (MVLI.ProcessedVarList.empty())
+    return nullptr;
+
+  return OMPUseDeviceAddrClause::Create(
+      Context, Locs, MVLI.ProcessedVarList, PrivateCopies, Inits,
+      MVLI.VarBaseDeclarations, MVLI.VarComponents);
+}
+
 OMPClause *Sema::ActOnOpenMPIsDevicePtrClause(ArrayRef<Expr *> VarList,
                                               const OMPVarListLocTy &Locs) {
   MappableVarListInfo MVLI(VarList);
   for (Expr *RefExpr : VarList) {
     assert(RefExpr && "NULL expr in OpenMP is_device_ptr clause.");
     SourceLocation ELoc;
     SourceRange ERange;
     Expr *SimpleRefExpr = RefExpr;

clang/lib/Sema/TreeTransform.h

   ///
   /// By default, performs semantic analysis to build the new OpenMP clause.
   /// Subclasses may override this routine to provide different behavior.
   OMPClause *RebuildOMPUseDevicePtrClause(ArrayRef<Expr *> VarList,
                                           const OMPVarListLocTy &Locs) {
     return getSema().ActOnOpenMPUseDevicePtrClause(VarList, Locs);
   }
 
+  /// Build a new OpenMP 'use_device_addr' clause.
+  ///
+  /// By default, performs semantic analysis to build the new OpenMP clause.
+  /// Subclasses may override this routine to provide different behavior.
+  OMPClause *RebuildOMPUseDeviceAddrClause(ArrayRef<Expr *> VarList,
+                                           const OMPVarListLocTy &Locs) {
+    return getSema().ActOnOpenMPUseDeviceAddrClause(VarList, Locs);
+  }
+
   /// Build a new OpenMP 'is_device_ptr' clause.
   ///
   /// By default, performs semantic analysis to build the new OpenMP clause.
   /// Subclasses may override this routine to provide different behavior.
   OMPClause *RebuildOMPIsDevicePtrClause(ArrayRef<Expr *> VarList,
                                          const OMPVarListLocTy &Locs) {
     return getSema().ActOnOpenMPIsDevicePtrClause(VarList, Locs);
   }
if (EVar.isInvalid())
       return nullptr;
     Vars.push_back(EVar.get());
   }
   OMPVarListLocTy Locs(C->getBeginLoc(), C->getLParenLoc(), C->getEndLoc());
   return getDerived().RebuildOMPUseDevicePtrClause(Vars, Locs);
 }
 
 template <typename Derived>
+OMPClause *TreeTransform<Derived>::TransformOMPUseDeviceAddrClause(
+    OMPUseDeviceAddrClause *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().RebuildOMPUseDeviceAddrClause(Vars, Locs);
+}
+
+template <typename Derived>
 OMPClause *
 TreeTransform<Derived>::TransformOMPIsDevicePtrClause(OMPIsDevicePtrClause *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;

clang/lib/Serialization/ASTReader.cpp

case llvm::omp::OMPC_use_device_ptr: {
     OMPMappableExprListSizeTy Sizes;
     Sizes.NumVars = Record.readInt();
     Sizes.NumUniqueDeclarations = Record.readInt();
     Sizes.NumComponentLists = Record.readInt();
     Sizes.NumComponents = Record.readInt();
     C = OMPUseDevicePtrClause::CreateEmpty(Context, Sizes);
     break;
   }
+  case llvm::omp::OMPC_use_device_addr: {
+    OMPMappableExprListSizeTy Sizes;
+    Sizes.NumVars = Record.readInt();
+    Sizes.NumUniqueDeclarations = Record.readInt();
+    Sizes.NumComponentLists = Record.readInt();
+    Sizes.NumComponents = Record.readInt();
+    C = OMPUseDeviceAddrClause::CreateEmpty(Context, Sizes);
+    break;
+  }
   case llvm::omp::OMPC_is_device_ptr: {
     OMPMappableExprListSizeTy Sizes;
     Sizes.NumVars = Record.readInt();
     Sizes.NumUniqueDeclarations = Record.readInt();
     Sizes.NumComponentLists = Record.readInt();
     Sizes.NumComponents = Record.readInt();
     C = OMPIsDevicePtrClause::CreateEmpty(Context, Sizes);
     break;
for (unsigned i = 0; i < TotalComponents; ++i) {
     Expr *AssociatedExpr = Record.readSubExpr();
     auto *AssociatedDecl = Record.readDeclAs<ValueDecl>();
     Components.push_back(OMPClauseMappableExprCommon::MappableComponent(
         AssociatedExpr, AssociatedDecl));
   }
   C->setComponents(Components, ListSizes);
 }
 
+void OMPClauseReader::VisitOMPUseDeviceAddrClause(OMPUseDeviceAddrClause *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();
+  for (unsigned i = 0; i != NumVars; ++i)
+    Vars.push_back(Record.readSubExpr());
+  C->setPrivateCopies(Vars);
+  Vars.clear();
+  for (unsigned i = 0; i != NumVars; ++i)
+    Vars.push_back(Record.readSubExpr());
+  C->setInits(Vars);
+
+  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.push_back(OMPClauseMappableExprCommon::MappableComponent(
+        AssociatedExpr, AssociatedDecl));
+  }
+  C->setComponents(Components, ListSizes);
+}
+
 void OMPClauseReader::VisitOMPIsDevicePtrClause(OMPIsDevicePtrClause *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;

clang/lib/Serialization/ASTWriter.cpp

void OMPClauseWriter::VisitOMPUseDevicePtrClause(OMPUseDevicePtrClause *C) {
   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::VisitOMPUseDeviceAddrClause(OMPUseDeviceAddrClause *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 *VE : C->private_copies())
+    Record.AddStmt(VE);
+  for (auto *VE : C->inits())
+    Record.AddStmt(VE);
+  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::VisitOMPIsDevicePtrClause(OMPIsDevicePtrClause *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);

clang/test/OpenMP/target_data_messages.c

-// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -ferror-limit 100 -o - %s -Wuninitialized
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp45 -fopenmp -fopenmp-version=45 -ferror-limit 100 -o - %s -Wuninitialized
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp50 -fopenmp -fopenmp-version=50 -ferror-limit 100 -o - %s -Wuninitialized
 
-// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp-simd -ferror-limit 100 -o - %s -Wuninitialized
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp45 -fopenmp-simd -fopenmp-version=45 -ferror-limit 100 -o - %s -Wuninitialized
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp50 -fopenmp-simd -fopenmp-version=50 -ferror-limit 100 -o - %s -Wuninitialized
 
 void foo() { }
 
 void xxx(int argc) {
   int map; // expected-note {{initialize the variable 'map' to silence this warning}}
 #pragma omp target data map(map) // expected-warning {{variable 'map' is uninitialized when used here}}
   for (int i = 0; i < 10; ++i)
     ;
 }
 
 int main(int argc, char **argv) {
   int a;
-  #pragma omp target data // expected-error {{expected at least one 'map' or 'use_device_ptr' clause for '#pragma omp target data'}}
+  #pragma omp target data // omp45-error {{expected at least one 'map' or 'use_device_ptr' clause for '#pragma omp target data'}} omp50-error {{expected at least one 'map', 'use_device_ptr', or 'use_device_addr' clause for '#pragma omp target data'}}
   {}
   L1:
     foo();
   #pragma omp target data map(a) allocate(a) // expected-error {{unexpected OpenMP clause 'allocate' in directive '#pragma omp target data'}}
   {
     foo();
     goto L1; // expected-error {{use of undeclared label 'L1'}}
   }

clang/test/OpenMP/target_data_use_device_ptr_addr_ast_print.cpp

-// RxUN: %clang_cc1 -verify -fopenmp -std=c++11 -ast-print %s | FileCheck %s
-// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -emit-pch -o %t %s
-// RUN: %clang_cc1 -fopenmp -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
-
-// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -emit-pch -o %t %s
-// RUN: %clang_cc1 -fopenmp-simd -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -std=c++11 -ast-print %s | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+
+// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=50 -std=c++11 -ast-print %s | FileCheck %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
 // expected-no-diagnostics
 
 #ifndef HEADER
 #define HEADER
 
 struct ST {
   int *a;
 };
 struct SA {
   int i, j;
   int *k = &j;
   int *&z = k;
+  int &y = i;
   void func(int arg) {
-#pragma omp target data map(tofrom: i) use_device_ptr(k)
+#pragma omp target data map(tofrom: i) use_device_ptr(k) use_device_addr(i, j)
     {}
-#pragma omp target data map(tofrom: i) use_device_ptr(z)
+#pragma omp target data map(tofrom: i) use_device_ptr(z) use_device_addr(k, y)
     {}
   return;
  }
 };
 // CHECK: struct SA
 // CHECK: void func(
-// CHECK: #pragma omp target data map(tofrom: this->i) use_device_ptr(this->k){{$}}
-// CHECK: #pragma omp target data map(tofrom: this->i) use_device_ptr(this->z)
+// CHECK: #pragma omp target data map(tofrom: this->i) use_device_ptr(this->k) use_device_addr(this->i,this->j){{$}}
+// CHECK: #pragma omp target data map(tofrom: this->i) use_device_ptr(this->z) use_device_addr(this->k,this->y)
 struct SB {
   unsigned A;
   unsigned B;
   float Arr[100];
   float *Ptr;
   float *foo() {
     return &Arr[0];
   }

 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 int from;
 
 template <typename T>
 T tmain(T argc) {
   T i;
   T &j = i;
   T *k = &j;
   T *&z = k;
 #pragma omp target data map(tofrom: i) use_device_ptr(k)
   {}
 #pragma omp target data map(tofrom: i) use_device_ptr(z)
   {}
   return 0;
 }
 
 // CHECK: template<> int tmain<int>(int argc) {
 // CHECK-NEXT: int i;
 // CHECK-NEXT: int &j = i;
 // CHECK-NEXT: int *k = &j;
int main(int argc, char **argv) {
   int i;
   int &j = i;
   int *k = &j;
   int *&z = k;
 // CHECK-NEXT: int i;
 // CHECK-NEXT: int &j = i;
 // CHECK-NEXT: int *k = &j;
 // CHECK-NEXT: int *&z = k;
-#pragma omp target data map(tofrom: i) use_device_ptr(k)
-// CHECK-NEXT: #pragma omp target data map(tofrom: i) use_device_ptr(k)
+#pragma omp target data map(tofrom: i) use_device_ptr(k) use_device_addr(i, j)
+// CHECK-NEXT: #pragma omp target data map(tofrom: i) use_device_ptr(k) use_device_addr(i,j)
   {}
 // CHECK-NEXT: {
 // CHECK-NEXT: }
-#pragma omp target data map(tofrom: i) use_device_ptr(z)
-// CHECK-NEXT: #pragma omp target data map(tofrom: i) use_device_ptr(z)
+#pragma omp target data map(tofrom: i) use_device_ptr(z) use_device_addr(i, j, k[:i])
+// CHECK-NEXT: #pragma omp target data map(tofrom: i) use_device_ptr(z) use_device_addr(i,j,k[:i])
   {}
   return tmain<int>(argc) + (*tmain<int*>(&argc));
 }
 
 #endif

clang/test/OpenMP/target_data_use_device_ptr_addr_messages.cpp

+// RUN: %clang_cc1 -std=c++11 -verify=expected,omp45 -fopenmp -fopenmp-version=45 -ferror-limit 200 %s -Wuninitialized
+// RUN: %clang_cc1 -std=c++11 -verify=expected,omp50 -fopenmp -fopenmp-version=50 -ferror-limit 200 %s -Wuninitialized
+
+// RUN: %clang_cc1 -std=c++11 -verify=expected,omp45 -fopenmp-simd -fopenmp-version=45 -ferror-limit 200 %s -Wuninitialized
+// RUN: %clang_cc1 -std=c++11 -verify=expected,omp50 -fopenmp-simd -fopenmp-version=50 -ferror-limit 200 %s -Wuninitialized
+struct ST {
+  int *a;
+};
+struct SA {
+  const int d = 5;
+  const int da[5] = { 0 };
+  ST e;
+  ST g[10];
+  int i;
+  int &j = i;
+  int *k = &j;
+  int *&z = k;
+  int aa[10];
+  void func(int arg) {
+#pragma omp target data map(i) use_device_ptr // expected-error {{expected '(' after 'use_device_ptr'}}
+    {}
+#pragma omp target data map(i) use_device_ptr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+    {}
+#pragma omp target data map(i) use_device_ptr() // expected-error {{expected expression}}
+    {}
+#pragma omp target data map(i) use_device_ptr(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
+    {}
+#pragma omp target data map(i) use_device_ptr(arg // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+    {}
+#pragma omp target data map(i) use_device_ptr(i) // expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+    {}
+#pragma omp target data map(i) use_device_ptr(j) // expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+    {}
+#pragma omp target data map(i) use_device_ptr(k) // OK
+    {}
+#pragma omp target data map(i) use_device_ptr(z) // OK
+    {}
+#pragma omp target data map(i) use_device_ptr(aa) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+    {}
+#pragma omp target data map(i) use_device_ptr(e) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+    {}
+#pragma omp target data map(i) use_device_ptr(g) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+    {}
+#pragma omp target data map(i) use_device_ptr(k,i,j) // expected-error2 {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+    {}
+#pragma omp target data map(i) use_device_ptr(d) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+    {}
+#pragma omp target data map(i) use_device_ptr(da) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+    {}
+#pragma omp target data map(i) use_device_addr // expected-error {{expected '(' after 'use_device_addr'}} omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+    {}
+#pragma omp target data map(i) use_device_addr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}} omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+    {}
+#pragma omp target data map(i) use_device_addr() // expected-error {{expected expression}} omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+    {}
+#pragma omp target data map(i) use_device_addr(alloc) // expected-error {{use of undeclared identifier 'alloc'}} omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+    {}
+#pragma omp target data map(i) use_device_addr(arg // expected-error {{expected ')'}} expected-note {{to match this '('}} omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+    {}
+#pragma omp target data map(i) use_device_addr(i) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+    {}
+#pragma omp target data map(i) use_device_addr(j) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+    {}
+#pragma omp target data map(i) use_device_addr(k) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+    {}
+#pragma omp target data map(i) use_device_addr(z) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+    {}
+#pragma omp target data map(i) use_device_addr(aa) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+    {}
+#pragma omp target data map(i) use_device_addr(e) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+    {}
+#pragma omp target data map(i) use_device_addr(g) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+    {}
+#pragma omp target data map(i) use_device_addr(k,i,j) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+    {}
+#pragma omp target data map(i) use_device_addr(d) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+    {}
+#pragma omp target data map(i) use_device_addr(da) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+    {}
+  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 int from;
+
+template <typename T, int I>
+T tmain(T argc) {
+  const T d = 5;
+  const T da[5] = { 0 };
+  S4 e(4);
+  S5 g(5);
+  T i;
+  T &j = i;
+  T *k = &j;
+  T *&z = k;
+  T aa[10];
+#pragma omp target data map(i) use_device_ptr // expected-error {{expected '(' after 'use_device_ptr'}}
+  {}
+#pragma omp target data map(i) use_device_ptr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+  {}
+#pragma omp target data map(i) use_device_ptr() // expected-error {{expected expression}}
+  {}
+#pragma omp target data map(i) use_device_ptr(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
+  {}
+#pragma omp target data map(i) use_device_ptr(argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(i) // expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(j) // expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(k) // OK
+  {}
+#pragma omp target data map(i) use_device_ptr(z) // OK
+  {}
+#pragma omp target data map(i) use_device_ptr(aa) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(e) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(g) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(k,i,j) // expected-error2 {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(d) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(da) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_addr // expected-error {{expected '(' after 'use_device_addr'}} omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+  {}
+#pragma omp target data map(i) use_device_addr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}} omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+  {}
+#pragma omp target data map(i) use_device_addr() // expected-error {{expected expression}} omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+  {}
+#pragma omp target data map(i) use_device_addr(alloc) // expected-error {{use of undeclared identifier 'alloc'}} omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+  {}
+#pragma omp target data map(i) use_device_addr(argc // expected-error {{expected ')'}} expected-note {{to match this '('}} omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+  {}
+#pragma omp target data map(i) use_device_addr(i) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+  {}
+#pragma omp target data map(i) use_device_addr(j) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+  {}
+#pragma omp target data map(i) use_device_addr(k) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+  {}
+#pragma omp target data map(i) use_device_addr(z) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+  {}
+#pragma omp target data map(i) use_device_addr(aa) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+  {}
+#pragma omp target data map(i) use_device_addr(e) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+  {}
+#pragma omp target data map(i) use_device_addr(g) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+  {}
+#pragma omp target data map(i) use_device_addr(k,i,j) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+  {}
+#pragma omp target data map(i) use_device_addr(d) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+  {}
+#pragma omp target data map(i) use_device_addr(da) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+  {}
+  return 0;
+}
+
+int main(int argc, char **argv) {
+  const int d = 5;
+  const int da[5] = { 0 };
+  S4 e(4);
+  S5 g(5);
+  int i;
+  int &j = i;
+  int *k = &j;
+  int *&z = k;
+  int aa[10];
+#pragma omp target data map(i) use_device_ptr // expected-error {{expected '(' after 'use_device_ptr'}}
+  {}
+#pragma omp target data map(i) use_device_ptr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+  {}
+#pragma omp target data map(i) use_device_ptr() // expected-error {{expected expression}}
+  {}
+#pragma omp target data map(i) use_device_ptr(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
+  {}
+#pragma omp target data map(i) use_device_ptr(argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(i) // expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(j) // expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(k) // OK
+  {}
+#pragma omp target data map(i) use_device_ptr(z) // OK
+  {}
+#pragma omp target data map(i) use_device_ptr(aa) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(e) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(g) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(k,i,j) // expected-error2 {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(d) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_ptr(da) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
+  {}
+#pragma omp target data map(i) use_device_addr // expected-error {{expected '(' after 'use_device_addr'}} omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+  {}
+#pragma omp target data map(i) use_device_addr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}} omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+  {}
+#pragma omp target data map(i) use_device_addr() // expected-error {{expected expression}} omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+  {}
+#pragma omp target data map(i) use_device_addr(alloc) // expected-error {{use of undeclared identifier 'alloc'}} omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+  {}
+#pragma omp target data map(i) use_device_addr(argc // expected-error {{expected ')'}} expected-note {{to match this '('}} omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+  {}
+#pragma omp target data map(i) use_device_addr(i) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+  {}
+#pragma omp target data map(i) use_device_addr(j) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+  {}
+#pragma omp target data map(i) use_device_addr(k) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+  {}
+#pragma omp target data map(i) use_device_addr(z) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+  {}
+#pragma omp target data map(i) use_device_addr(aa) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+  {}
+#pragma omp target data map(i) use_device_addr(e) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+  {}
+#pragma omp target data map(i) use_device_addr(g) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+  {}
+#pragma omp target data map(i) use_device_addr(k,i,j) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+  {}
+#pragma omp target data map(i) use_device_addr(d) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+  {}
+#pragma omp target data map(i) use_device_addr(da) // omp45-error {{unexpected OpenMP clause 'use_device_addr' in directive '#pragma omp target data'}}
+  {}
+  return tmain<int, 3>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 3>' requested here}}
+}

clang/test/OpenMP/target_data_use_device_ptr_messages.cpp

-// RUN: %clang_cc1 -std=c++11 -verify -fopenmp -ferror-limit 200 %s -Wuninitialized
-
-// RUN: %clang_cc1 -std=c++11 -verify -fopenmp-simd -ferror-limit 200 %s -Wuninitialized
-struct ST {
-  int *a;
-};
-struct SA {
-  const int d = 5;
-  const int da[5] = { 0 };
-  ST e;
-  ST g[10];
-  int i;
-  int &j = i;
-  int *k = &j;
-  int *&z = k;
-  int aa[10];
-  void func(int arg) {
-#pragma omp target data map(i) use_device_ptr // expected-error {{expected '(' after 'use_device_ptr'}}
-    {}
-#pragma omp target data map(i) use_device_ptr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
-    {}
-#pragma omp target data map(i) use_device_ptr() // expected-error {{expected expression}}
-    {}
-#pragma omp target data map(i) use_device_ptr(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
-    {}
-#pragma omp target data map(i) use_device_ptr(arg // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
-    {}
-#pragma omp target data map(i) use_device_ptr(i) // expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
-    {}
-#pragma omp target data map(i) use_device_ptr(j) // expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
-    {}
-#pragma omp target data map(i) use_device_ptr(k) // OK
-    {}
-#pragma omp target data map(i) use_device_ptr(z) // OK
-    {}
-#pragma omp target data map(i) use_device_ptr(aa) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
-    {}
-#pragma omp target data map(i) use_device_ptr(e) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
-    {}
-#pragma omp target data map(i) use_device_ptr(g) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
-    {}
-#pragma omp target data map(i) use_device_ptr(k,i,j) // expected-error2 {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
-    {}
-#pragma omp target data map(i) use_device_ptr(d) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
-    {}
-#pragma omp target data map(i) use_device_ptr(da) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
-    {}
-  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 int from;
-
-template <typename T, int I>
-T tmain(T argc) {
-  const T d = 5;
-  const T da[5] = { 0 };
-  S4 e(4);
-  S5 g(5);
-  T i;
-  T &j = i;
-  T *k = &j;
-  T *&z = k;
-  T aa[10];
-#pragma omp target data map(i) use_device_ptr // expected-error {{expected '(' after 'use_device_ptr'}}
-  {}
-#pragma omp target data map(i) use_device_ptr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
-  {}
-#pragma omp target data map(i) use_device_ptr() // expected-error {{expected expression}}
-  {}
-#pragma omp target data map(i) use_device_ptr(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
-  {}
-#pragma omp target data map(i) use_device_ptr(argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
-  {}
-#pragma omp target data map(i) use_device_ptr(i) // expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
-  {}
-#pragma omp target data map(i) use_device_ptr(j) // expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
-  {}
-#pragma omp target data map(i) use_device_ptr(k) // OK
-  {}
-#pragma omp target data map(i) use_device_ptr(z) // OK
-  {}
-#pragma omp target data map(i) use_device_ptr(aa) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
-  {}
-#pragma omp target data map(i) use_device_ptr(e) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
-  {}
-#pragma omp target data map(i) use_device_ptr(g) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
-  {}
-#pragma omp target data map(i) use_device_ptr(k,i,j) // expected-error2 {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
-  {}
-#pragma omp target data map(i) use_device_ptr(d) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
-  {}
-#pragma omp target data map(i) use_device_ptr(da) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
-  {}
-  return 0;
-}
-
-int main(int argc, char **argv) {
-  const int d = 5;
-  const int da[5] = { 0 };
-  S4 e(4);
-  S5 g(5);
-  int i;
-  int &j = i;
-  int *k = &j;
-  int *&z = k;
-  int aa[10];
-#pragma omp target data map(i) use_device_ptr // expected-error {{expected '(' after 'use_device_ptr'}}
-  {}
-#pragma omp target data map(i) use_device_ptr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
-  {}
-#pragma omp target data map(i) use_device_ptr() // expected-error {{expected expression}}
-  {}
-#pragma omp target data map(i) use_device_ptr(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
-  {}
-#pragma omp target data map(i) use_device_ptr(argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
-  {}
-#pragma omp target data map(i) use_device_ptr(i) // expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
-  {}
-#pragma omp target data map(i) use_device_ptr(j) // expected-error {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
-  {}
-#pragma omp target data map(i) use_device_ptr(k) // OK
-  {}
-#pragma omp target data map(i) use_device_ptr(z) // OK
-  {}
-#pragma omp target data map(i) use_device_ptr(aa) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
-  {}
-#pragma omp target data map(i) use_device_ptr(e) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
-  {}
-#pragma omp target data map(i) use_device_ptr(g) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
-  {}
-#pragma omp target data map(i) use_device_ptr(k,i,j) // expected-error2 {{expected pointer or reference to pointer in 'use_device_ptr' clause}}
-  {}
-#pragma omp target data map(i) use_device_ptr(d) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
-  {}
-#pragma omp target data map(i) use_device_ptr(da) // expected-error{{expected pointer or reference to pointer in 'use_device_ptr' clause}}
-  {}
-  return tmain<int, 3>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 3>' requested here}}
-}

clang/test/OpenMP/target_map_messages.cpp

int main(int argc, char **argv) {
   int *k = &j;
   S6<int> m;
   int x;
   int y;
   int to, tofrom, always, close;
   const int (&l)[5] = da;
   SC1 s;
   SC1 *p;
-#pragma omp target data map // expected-error {{expected '(' after 'map'}} expected-error {{expected at least one 'map' or 'use_device_ptr' clause for '#pragma omp target data'}}
+#pragma omp target data map // expected-error {{expected '(' after 'map'}} le45-error {{expected at least one 'map' or 'use_device_ptr' clause for '#pragma omp target data'}} le50-error {{expected at least one 'map', 'use_device_ptr', or 'use_device_addr' clause for '#pragma omp target data'}}
 #pragma omp target data map( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
 #pragma omp target data map() // expected-error {{expected expression}}
 #pragma omp target data map(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
 #pragma omp target data map(to argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected ',' or ')' in 'map' clause}}
 #pragma omp target data map(to:) // expected-error {{expected expression}}
 #pragma omp target data map(from: argc, // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
 #pragma omp target data map(x: y) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
 #pragma omp target map(x)

clang/test/OpenMP/target_teams_map_messages.cpp

int main(int argc, char **argv) {
   int i;
   int &j = i;
   int *k = &j;
   S6<int> m;
   int x;
   int y;
   int to, tofrom, always;
   const int (&l)[5] = da;
-#pragma omp target data map // expected-error {{expected '(' after 'map'}} expected-error {{expected at least one 'map' or 'use_device_ptr' clause for '#pragma omp target data'}}
+#pragma omp target data map // expected-error {{expected '(' after 'map'}} le45-error {{expected at least one 'map' or 'use_device_ptr' clause for '#pragma omp target data'}} le50-error {{expected at least one 'map', 'use_device_ptr', or 'use_device_addr' clause for '#pragma omp target data'}}
 #pragma omp target data map( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
 #pragma omp target data map() // expected-error {{expected expression}}
 #pragma omp target data map(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
 #pragma omp target data map(to argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected ',' or ')' in 'map' clause}}
 #pragma omp target data map(to:) // expected-error {{expected expression}}
 #pragma omp target data map(from: argc, // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
 #pragma omp target data map(x: y) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
 #pragma omp target teams map(x)

clang/tools/libclang/CIndex.cpp

 }
 void OMPClauseEnqueue::VisitOMPFromClause(const OMPFromClause *C) {
   VisitOMPClauseList(C);
 }
 void OMPClauseEnqueue::VisitOMPUseDevicePtrClause(
     const OMPUseDevicePtrClause *C) {
   VisitOMPClauseList(C);
 }
+void OMPClauseEnqueue::VisitOMPUseDeviceAddrClause(
+    const OMPUseDeviceAddrClause *C) {
+  VisitOMPClauseList(C);
+}
 void OMPClauseEnqueue::VisitOMPIsDevicePtrClause(
     const OMPIsDevicePtrClause *C) {
   VisitOMPClauseList(C);
 }
 void OMPClauseEnqueue::VisitOMPNontemporalClause(
     const OMPNontemporalClause *C) {
   VisitOMPClauseList(C);
   for (const auto *E : C->private_refs())

llvm/include/llvm/Frontend/OpenMP/OMPKinds.def

 __OMP_CLAUSE(nontemporal, OMPNontemporalClause)
 __OMP_CLAUSE(order, OMPOrderClause)
 __OMP_CLAUSE(destroy, OMPDestroyClause)
 __OMP_CLAUSE(detach, OMPDetachClause)
 __OMP_CLAUSE(inclusive, OMPInclusiveClause)
 __OMP_CLAUSE(exclusive, OMPExclusiveClause)
 __OMP_CLAUSE(uses_allocators, OMPUsesAllocatorsClause)
 __OMP_CLAUSE(affinity, OMPAffinityClause)
+__OMP_CLAUSE(use_device_addr, OMPUseDeviceAddrClause)
 
 __OMP_CLAUSE_NO_CLASS(uniform)
 __OMP_CLAUSE_NO_CLASS(device_type)
 __OMP_CLAUSE_NO_CLASS(match)
 
 __OMP_IMPLICIT_CLAUSE_CLASS(depobj, "depobj", OMPDepobjClause)
 __OMP_IMPLICIT_CLAUSE_CLASS(flush, "flush", OMPFlushClause)
 

 __OMP_DIRECTIVE_CLAUSE(requires, 1, ~0, atomic_default_mem_order)
 
 __OMP_DIRECTIVE_CLAUSE(allocate, 1, ~0, allocator)
 
 __OMP_DIRECTIVE_CLAUSE(target_data, 1, ~0, if)
 __OMP_DIRECTIVE_CLAUSE(target_data, 1, ~0, device)
 __OMP_DIRECTIVE_CLAUSE(target_data, 1, ~0, map)
 __OMP_DIRECTIVE_CLAUSE(target_data, 1, ~0, use_device_ptr)
+__OMP_DIRECTIVE_CLAUSE(target_data, 50, ~0, use_device_addr)
 
 __OMP_DIRECTIVE_CLAUSE(target_enter_data, 1, ~0, if)
 __OMP_DIRECTIVE_CLAUSE(target_enter_data, 1, ~0, device)
 __OMP_DIRECTIVE_CLAUSE(target_enter_data, 1, ~0, map)
 __OMP_DIRECTIVE_CLAUSE(target_enter_data, 1, ~0, nowait)
 __OMP_DIRECTIVE_CLAUSE(target_enter_data, 1, ~0, depend)
 
 __OMP_DIRECTIVE_CLAUSE(target_exit_data, 1, ~0, if)