[clang][OpenMP5.1] Initial parsing/sema for has_device_addr

[jyu2]

Summary

Added basic parsing/sema/ support for the 'has_device_addr' clause

[ABataev]

I think it worth it to add the analysis test for the clause

clang/lib/AST/OpenMPClause.cpp

1465

auto *

clang/lib/Serialization/ASTReader.cpp

12848

Use I instead of i everywhere in this function

llvm/include/llvm/Frontend/OpenMP/OMP.td

564

VersionedClause<OMPC_HasDeviceAddr, 51>,

661

VersionedClause<OMPC_HasDeviceAddr, 51>,

694

VersionedClause<OMPC_HasDeviceAddr, 51>,

711

VersionedClause<OMPC_HasDeviceAddr, 51>,

1145

VersionedClause<OMPC_HasDeviceAddr, 51>,

1176

VersionedClause<OMPC_HasDeviceAddr, 51>, and in other places

[jyu2]

Thanks Alexey, for the review! This patch addresses his comments.

[ABataev]

Add a test to Analysis/cfg-openmp.cpp

[jyu2]

Add test in cfg-openmp.cpp

[ABataev]

LG

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_is_device_ptr;
   }
 };
 
+/// 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
 /// #pragma omp simd nontemporal(a)
 /// \endcode
 /// In this example directive '#pragma omp simd' has clause 'nontemporal' for
 /// the variable 'a'.
 class OMPNontemporalClause final

clang/include/clang/AST/RecursiveASTVisitor.h

 template <typename Derived>
 bool RecursiveASTVisitor<Derived>::VisitOMPIsDevicePtrClause(
     OMPIsDevicePtrClause *C) {
   TRY_TO(VisitOMPClauseList(C));
   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) {
   TRY_TO(VisitOMPClauseList(C));
   for (auto *E : C->private_refs()) {
     TRY_TO(TraverseStmt(E));
   }
   return true;
 }

clang/include/clang/Sema/Sema.h

   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 '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,
                                           SourceLocation LParenLoc,
                                           SourceLocation EndLoc);
 
   /// Data for list of allocators.
   struct UsesAllocatorsData {

clang/lib/AST/OpenMPClause.cpp

const OMPClauseWithPreInit *OMPClauseWithPreInit::get(const OMPClause *C) {
   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_has_device_addr:
   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:
   case OMPC_nontemporal:
const OMPClauseWithPostUpdate *OMPClauseWithPostUpdate::get(const OMPClause *C) {
   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_has_device_addr:
   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:
   case OMPC_nontemporal:
void *Mem = C.Allocate(
       totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
                        OMPClauseMappableExprCommon::MappableComponent>(
           Sizes.NumVars, Sizes.NumUniqueDeclarations,
           Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
           Sizes.NumComponents));
   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,
                                                    SourceLocation EndLoc,
                                                    ArrayRef<Expr *> VL) {
   // Allocate space for nontemporal variables + private references.
   void *Mem = C.Allocate(totalSizeToAlloc<Expr *>(2 * VL.size()));
   auto *Clause =

 void OMPClausePrinter::VisitOMPIsDevicePtrClause(OMPIsDevicePtrClause *Node) {
   if (!Node->varlist_empty()) {
     OS << "is_device_ptr";
     VisitOMPClauseList(Node, '(');
     OS << ")";
   }
 }
 
+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";
     VisitOMPClauseList(Node, '(');
     OS << ")";
   }
 }
 

clang/lib/AST/StmtProfile.cpp

 void OMPClauseProfiler::VisitOMPUseDeviceAddrClause(
     const OMPUseDeviceAddrClause *C) {
   VisitOMPClauseList(C);
 }
 void OMPClauseProfiler::VisitOMPIsDevicePtrClause(
     const OMPIsDevicePtrClause *C) {
   VisitOMPClauseList(C);
 }
+void OMPClauseProfiler::VisitOMPHasDeviceAddrClause(
+    const OMPHasDeviceAddrClause *C) {
+  VisitOMPClauseList(C);
+}
 void OMPClauseProfiler::VisitOMPNontemporalClause(
     const OMPNontemporalClause *C) {
   VisitOMPClauseList(C);
   for (auto *E : C->private_refs())
     Profiler->VisitStmt(E);
 }
 void OMPClauseProfiler::VisitOMPInclusiveClause(const OMPInclusiveClause *C) {
   VisitOMPClauseList(C);

clang/lib/Basic/OpenMPKinds.cpp

#include "clang/Basic/OpenMPKinds.def"
   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_has_device_addr:
   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:
   case OMPC_novariants:
#include "clang/Basic/OpenMPKinds.def"
   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_has_device_addr:
   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:
   case OMPC_detach:

clang/lib/CodeGen/CGStmtOpenMP.cpp

static void emitOMPAtomicExpr(CodeGenFunction &CGF, OpenMPClauseKind Kind,
   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_has_device_addr:
   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:
   case OMPC_nontemporal:

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 | 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) {
   OMPClauseKind = CKind;
   OMPClause *Clause = nullptr;
   bool ErrorFound = false;
   bool WrongDirective = false;
   // Check if clause is allowed for the given directive.
OMPClause *Parser::ParseOpenMPClause(OpenMPDirectiveKind DKind,
   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_has_device_addr:
   case OMPC_allocate:
   case OMPC_nontemporal:
   case OMPC_inclusive:
   case OMPC_exclusive:
   case OMPC_affinity:
     Clause = ParseOpenMPVarListClause(DKind, CKind, WrongDirective);
     break;
   case OMPC_sizes:

 ///    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 ')'
+///    has_device_addr-clause:
+///       'has_device_addr' '(' list ')'
 ///    allocate-clause:
 ///       'allocate' '(' [ allocator ':' ] list ')'
 ///    nontemporal-clause:
 ///       'nontemporal' '(' list ')'
 ///    inclusive-clause:
 ///       'inclusive' '(' list ')'
 ///    exclusive-clause:
 ///       'exclusive' '(' list ')'

clang/lib/Sema/SemaOpenMP.cpp

for (OMPClause *C : Clauses) {
       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_has_device_addr:
       case OMPC_nontemporal:
       case OMPC_order:
       case OMPC_destroy:
       case OMPC_inclusive:
       case OMPC_exclusive:
       case OMPC_uses_allocators:
       case OMPC_affinity:
       case OMPC_bind:
OMPClause *Sema::ActOnOpenMPSimpleClause(
   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_has_device_addr:
   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:
   case OMPC_destroy:
OMPClause *Sema::ActOnOpenMPSingleExprWithArgClause(
   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_has_device_addr:
   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:
   case OMPC_nontemporal:
OMPClause *Sema::ActOnOpenMPClause(OpenMPClauseKind Kind,
   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_has_device_addr:
   case OMPC_atomic_default_mem_order:
   case OMPC_device_type:
   case OMPC_match:
   case OMPC_nontemporal:
   case OMPC_order:
   case OMPC_novariants:
   case OMPC_nocontext:
   case OMPC_detach:
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_has_device_addr:
+    Res = ActOnOpenMPHasDeviceAddrClause(VarList, Locs);
+    break;
   case OMPC_allocate:
     Res = ActOnOpenMPAllocateClause(DepModOrTailExpr, VarList, StartLoc,
                                     LParenLoc, ColonLoc, EndLoc);
     break;
   case OMPC_nontemporal:
     Res = ActOnOpenMPNontemporalClause(VarList, StartLoc, LParenLoc, EndLoc);
     break;
   case OMPC_inclusive:
OMPClause *Sema::ActOnOpenMPIsDevicePtrClause(ArrayRef<Expr *> VarList,
   if (MVLI.ProcessedVarList.empty())
     return nullptr;
 
   return OMPIsDevicePtrClause::Create(Context, Locs, MVLI.ProcessedVarList,
                                       MVLI.VarBaseDeclarations,
                                       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) {
   if (Allocator) {
     // OpenMP [2.11.4 allocate Clause, Description]
     // allocator is an expression of omp_allocator_handle_t type.
     if (!findOMPAllocatorHandleT(*this, Allocator->getExprLoc(), DSAStack))
       return nullptr;

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 *RebuildOMPIsDevicePtrClause(ArrayRef<Expr *> VarList,
                                          const OMPVarListLocTy &Locs) {
     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.
   /// Subclasses may override this routine to provide different behavior.
   OMPClause *RebuildOMPDefaultmapClause(OpenMPDefaultmapClauseModifier M,
                                         OpenMPDefaultmapClauseKind Kind,
                                         SourceLocation StartLoc,
                                         SourceLocation LParenLoc,
if (EVar.isInvalid())
       return nullptr;
     Vars.push_back(EVar.get());
   }
   OMPVarListLocTy Locs(C->getBeginLoc(), C->getLParenLoc(), C->getEndLoc());
   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) {
   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_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;
   }
+  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;
   case llvm::omp::OMPC_nontemporal:
     C = OMPNontemporalClause::CreateEmpty(Context, Record.readInt());
     break;
   case llvm::omp::OMPC_inclusive:
     C = OMPInclusiveClause::CreateEmpty(Context, Record.readInt());
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::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();
   SmallVector<Expr *, 16> Vars;
   Vars.reserve(NumVars);
   for (unsigned i = 0; i != NumVars; ++i)
     Vars.push_back(Record.readSubExpr());
   C->setVarRefs(Vars);

clang/lib/Serialization/ASTWriter.cpp

void OMPClauseWriter::VisitOMPIsDevicePtrClause(OMPIsDevicePtrClause *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::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(
     OMPUnifiedSharedMemoryClause *) {}
 
 void OMPClauseWriter::VisitOMPReverseOffloadClause(OMPReverseOffloadClause *) {}
 
 void

clang/test/Analysis/cfg-openmp.cpp

 // 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]
 // CHECK-NEXT:   1: int x;
 // CHECK-NEXT:   2: int cond;
 // CHECK-NEXT:   3: int fp;
 // CHECK-NEXT:   4: int rd;
 // CHECK-NEXT:   5: int lin;

 // CHECK-DAG:  [[#TTD+6]]: [B3.[[#TTDB]]]
 // CHECK-DAG:  [[#TTD+7]]: #pragma omp target parallel loop if(cond) firstprivate(fp) reduction(+: rd) map(alloc: map)
 // CHECK-DAG:    for (int i = 0;
 // CHECK-DAG:        [B3.[[#TTDB+3]]];
 #pragma omp target parallel loop if(cond) firstprivate(fp) reduction(+:rd) map(alloc:map)
   for (int i = 0; i < 10; ++i)
     argc = x;
 }
+
+#endif  // _OPENMP

clang/test/OpenMP/target_has_device_addr_ast_print.cpp

+// 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

clang/test/OpenMP/target_has_device_addr_messages.cpp

+// 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);
+}

clang/tools/libclang/CIndex.cpp

 void OMPClauseEnqueue::VisitOMPUseDeviceAddrClause(
     const OMPUseDeviceAddrClause *C) {
   VisitOMPClauseList(C);
 }
 void OMPClauseEnqueue::VisitOMPIsDevicePtrClause(
     const OMPIsDevicePtrClause *C) {
   VisitOMPClauseList(C);
 }
+void OMPClauseEnqueue::VisitOMPHasDeviceAddrClause(
+    const OMPHasDeviceAddrClause *C) {
+  VisitOMPClauseList(C);
+}
 void OMPClauseEnqueue::VisitOMPNontemporalClause(
     const OMPNontemporalClause *C) {
   VisitOMPClauseList(C);
   for (const auto *E : C->private_refs())
     Visitor->AddStmt(E);
 }
 void OMPClauseEnqueue::VisitOMPOrderClause(const OMPOrderClause *C) {}
 void OMPClauseEnqueue::VisitOMPUsesAllocatorsClause(

flang/lib/Semantics/check-omp-structure.cpp

 CHECK_SIMPLE_CLAUSE(Nontemporal, OMPC_nontemporal)
 CHECK_SIMPLE_CLAUSE(Order, OMPC_order)
 CHECK_SIMPLE_CLAUSE(Read, OMPC_read)
 CHECK_SIMPLE_CLAUSE(ReverseOffload, OMPC_reverse_offload)
 CHECK_SIMPLE_CLAUSE(Threadprivate, OMPC_threadprivate)
 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)
 CHECK_SIMPLE_CLAUSE(Nogroup, OMPC_nogroup)
 CHECK_SIMPLE_CLAUSE(Notinbranch, OMPC_notinbranch)
 CHECK_SIMPLE_CLAUSE(Nowait, OMPC_nowait)
 CHECK_SIMPLE_CLAUSE(Partial, OMPC_partial)
 CHECK_SIMPLE_CLAUSE(ProcBind, OMPC_proc_bind)

llvm/include/llvm/Frontend/OpenMP/OMP.td

def OMPC_UseDevicePtr : Clause<"use_device_ptr"> {
   let flangClass = "Name";
   let isValueList = true;
 }
 def OMPC_IsDevicePtr : Clause<"is_device_ptr"> {
   let clangClass = "OMPIsDevicePtrClause";
   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";
 }
 def OMPC_InReduction : Clause<"in_reduction"> {
   let clangClass = "OMPInReductionClause";
 }
 def OMPC_UnifiedAddress : Clause<"unified_address"> {

 def OMP_Target : Directive<"target"> {
   let allowedClauses = [
     VersionedClause<OMPC_If>,
     VersionedClause<OMPC_Map>,
     VersionedClause<OMPC_Private>,
     VersionedClause<OMPC_Depend>,
     VersionedClause<OMPC_FirstPrivate>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_Reduction>,
     VersionedClause<OMPC_Allocate>,
     VersionedClause<OMPC_UsesAllocators, 50>
   ];
   let allowedOnceClauses = [
     VersionedClause<OMPC_Device>,
     VersionedClause<OMPC_DefaultMap>,
     VersionedClause<OMPC_NoWait>
let allowedClauses = [
     VersionedClause<OMPC_NoWait>,
     VersionedClause<OMPC_Depend>,
     VersionedClause<OMPC_Private>,
     VersionedClause<OMPC_FirstPrivate>,
     VersionedClause<OMPC_Default>,
     VersionedClause<OMPC_Shared>,
     VersionedClause<OMPC_Reduction>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_Allocate>,
     VersionedClause<OMPC_UsesAllocators, 50>
   ];
   let allowedOnceClauses = [
     VersionedClause<OMPC_DefaultMap>,
     VersionedClause<OMPC_Device>,
     VersionedClause<OMPC_If>,
     VersionedClause<OMPC_NumThreads>,
let allowedClauses = [
     VersionedClause<OMPC_ProcBind>,
     VersionedClause<OMPC_Shared>,
     VersionedClause<OMPC_Reduction>,
     VersionedClause<OMPC_Collapse>,
     VersionedClause<OMPC_Schedule>,
     VersionedClause<OMPC_Ordered>,
     VersionedClause<OMPC_Linear>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_Allocate>,
     VersionedClause<OMPC_Order, 50>,
     VersionedClause<OMPC_UsesAllocators, 50>
   ];
 }
 def OMP_TargetParallelDo : Directive<"target parallel do"> {
   let allowedClauses = [
     VersionedClause<OMPC_Map>,
     VersionedClause<OMPC_Private>,
     VersionedClause<OMPC_FirstPrivate>,
     VersionedClause<OMPC_LastPrivate>,
     VersionedClause<OMPC_Depend>,
     VersionedClause<OMPC_Shared>,
     VersionedClause<OMPC_Reduction>,
     VersionedClause<OMPC_Linear>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_Allocator>,
     VersionedClause<OMPC_Order>,
     VersionedClause<OMPC_UsesAllocators>,
     VersionedClause<OMPC_Default>,
     VersionedClause<OMPC_Copyin>
   ];
   let allowedOnceClauses = [
     VersionedClause<OMPC_If>,
let allowedClauses = [
     VersionedClause<OMPC_Collapse>,
     VersionedClause<OMPC_Schedule>,
     VersionedClause<OMPC_Ordered>,
     VersionedClause<OMPC_Linear>,
     VersionedClause<OMPC_SafeLen>,
     VersionedClause<OMPC_SimdLen>,
     VersionedClause<OMPC_Aligned>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_Allocate>,
     VersionedClause<OMPC_NonTemporal, 50>,
     VersionedClause<OMPC_Order, 50>,
     VersionedClause<OMPC_UsesAllocators, 50>
   ];
 }
 def OMP_TargetParallelDoSimd : Directive<"target parallel do simd"> {
   let allowedClauses = [
let allowedClauses = [
     VersionedClause<OMPC_Collapse>,
     VersionedClause<OMPC_Schedule>,
     VersionedClause<OMPC_Ordered>,
     VersionedClause<OMPC_Linear>,
     VersionedClause<OMPC_SafeLen>,
     VersionedClause<OMPC_SimdLen>,
     VersionedClause<OMPC_Aligned>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_Allocate>,
     VersionedClause<OMPC_NonTemporal>,
     VersionedClause<OMPC_Order>,
     VersionedClause<OMPC_UsesAllocators>
   ];
 }
 def OMP_TargetSimd : Directive<"target simd"> {
   let allowedClauses = [
     VersionedClause<OMPC_Aligned>,
     VersionedClause<OMPC_Allocate>,
     VersionedClause<OMPC_Depend>,
     VersionedClause<OMPC_FirstPrivate>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_LastPrivate>,
     VersionedClause<OMPC_Linear>,
     VersionedClause<OMPC_Map>,
     VersionedClause<OMPC_NonTemporal, 50>,
     VersionedClause<OMPC_NoWait>,
     VersionedClause<OMPC_Order, 50>,
     VersionedClause<OMPC_Private>,
     VersionedClause<OMPC_Reduction>,

 def OMP_TargetTeams : Directive<"target teams"> {
   let allowedClauses = [
     VersionedClause<OMPC_If>,
     VersionedClause<OMPC_Map>,
     VersionedClause<OMPC_Private>,
     VersionedClause<OMPC_Depend>,
     VersionedClause<OMPC_FirstPrivate>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_Reduction>,
     VersionedClause<OMPC_Allocate>,
     VersionedClause<OMPC_UsesAllocators, 50>,
     VersionedClause<OMPC_Shared>
   ];
 
   let allowedOnceClauses = [
     VersionedClause<OMPC_Device>,
     VersionedClause<OMPC_NoWait>,
     VersionedClause<OMPC_DefaultMap>,
     VersionedClause<OMPC_Default>,
     VersionedClause<OMPC_NumTeams>,
     VersionedClause<OMPC_ThreadLimit>
   ];
 }
 def OMP_TargetTeamsDistribute : Directive<"target teams distribute"> {
   let allowedClauses = [
     VersionedClause<OMPC_If>,
     VersionedClause<OMPC_Map>,
     VersionedClause<OMPC_Private>,
     VersionedClause<OMPC_Depend>,
     VersionedClause<OMPC_FirstPrivate>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_Reduction>,
     VersionedClause<OMPC_Allocate>,
     VersionedClause<OMPC_UsesAllocators, 50>,
     VersionedClause<OMPC_Shared>,
     VersionedClause<OMPC_LastPrivate>
   ];
   let allowedOnceClauses = [
     VersionedClause<OMPC_Device>,
let allowedClauses = [
     VersionedClause<OMPC_Device>,
     VersionedClause<OMPC_Map>,
     VersionedClause<OMPC_Private>,
     VersionedClause<OMPC_NoWait>,
     VersionedClause<OMPC_Depend>,
     VersionedClause<OMPC_DefaultMap>,
     VersionedClause<OMPC_FirstPrivate>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_Default>,
     VersionedClause<OMPC_Shared>,
     VersionedClause<OMPC_Reduction>,
     VersionedClause<OMPC_NumTeams>,
     VersionedClause<OMPC_ThreadLimit>,
     VersionedClause<OMPC_LastPrivate>,
     VersionedClause<OMPC_Collapse>,
     VersionedClause<OMPC_DistSchedule>,
def OMP_TargetTeamsDistributeParallelDo :
     Directive<"target teams distribute parallel do"> {
   let allowedClauses = [
     VersionedClause<OMPC_If>,
     VersionedClause<OMPC_Map>,
     VersionedClause<OMPC_Private>,
     VersionedClause<OMPC_Depend>,
     VersionedClause<OMPC_FirstPrivate>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_Reduction>,
     VersionedClause<OMPC_Allocate>,
     VersionedClause<OMPC_UsesAllocators>,
     VersionedClause<OMPC_Shared>,
     VersionedClause<OMPC_LastPrivate>,
     VersionedClause<OMPC_Copyin>,
     VersionedClause<OMPC_Linear>,
     VersionedClause<OMPC_Ordered>,
let allowedClauses = [
     VersionedClause<OMPC_Device>,
     VersionedClause<OMPC_Map>,
     VersionedClause<OMPC_Private>,
     VersionedClause<OMPC_NoWait>,
     VersionedClause<OMPC_Depend>,
     VersionedClause<OMPC_DefaultMap>,
     VersionedClause<OMPC_FirstPrivate>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_Default>,
     VersionedClause<OMPC_Shared>,
     VersionedClause<OMPC_Reduction>,
     VersionedClause<OMPC_NumTeams>,
     VersionedClause<OMPC_ThreadLimit>,
     VersionedClause<OMPC_LastPrivate>,
     VersionedClause<OMPC_Collapse>,
     VersionedClause<OMPC_DistSchedule>,

 def OMP_TargetTeamsDistributeParallelDoSimd :
     Directive<"target teams distribute parallel do simd"> {
   let allowedClauses = [
     VersionedClause<OMPC_Map>,
     VersionedClause<OMPC_Private>,
     VersionedClause<OMPC_Depend>,
     VersionedClause<OMPC_FirstPrivate>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_Reduction>,
     VersionedClause<OMPC_Allocate>,
     VersionedClause<OMPC_UsesAllocators>,
     VersionedClause<OMPC_Shared>,
     VersionedClause<OMPC_LastPrivate>,
     VersionedClause<OMPC_Copyin>,
     VersionedClause<OMPC_Linear>,
     VersionedClause<OMPC_Ordered>,
def OMP_TargetTeamsDistributeSimd :
     Directive<"target teams distribute simd"> {
   let allowedClauses = [
     VersionedClause<OMPC_Aligned>,
     VersionedClause<OMPC_Allocate>,
     VersionedClause<OMPC_Depend>,
     VersionedClause<OMPC_FirstPrivate>,
     VersionedClause<OMPC_If>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_LastPrivate>,
     VersionedClause<OMPC_Linear>,
     VersionedClause<OMPC_Map>,
     VersionedClause<OMPC_NonTemporal, 50>,
     VersionedClause<OMPC_Order, 50>,
     VersionedClause<OMPC_Private>,
     VersionedClause<OMPC_Reduction>,
     VersionedClause<OMPC_Shared>,
let allowedClauses = [
     VersionedClause<OMPC_NoWait>,
     VersionedClause<OMPC_Use>,
   ];
 }
 def OMP_dispatch : Directive<"dispatch"> {
   let allowedClauses = [
     VersionedClause<OMPC_Device>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_NoWait>,
     VersionedClause<OMPC_Depend>,
     VersionedClause<OMPC_Novariants>,
     VersionedClause<OMPC_Nocontext>
   ];
 }
 def OMP_masked : Directive<"masked"> {
   let allowedOnceClauses = [

 def OMP_target_teams_loop : Directive<"target teams loop"> {
   let allowedClauses = [
     VersionedClause<OMPC_Allocate>,
     VersionedClause<OMPC_Depend>,
     VersionedClause<OMPC_DefaultMap>,
     VersionedClause<OMPC_Device>,
     VersionedClause<OMPC_FirstPrivate>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_LastPrivate>,
     VersionedClause<OMPC_Map>,
     VersionedClause<OMPC_Private>,
     VersionedClause<OMPC_Reduction>,
     VersionedClause<OMPC_Shared>,
     VersionedClause<OMPC_UsesAllocators, 50>
   ];
   let allowedOnceClauses = [

 def OMP_target_parallel_loop : Directive<"target parallel loop"> {
   let allowedClauses = [
     VersionedClause<OMPC_Allocate>,
     VersionedClause<OMPC_Copyin>,
     VersionedClause<OMPC_Depend>,
     VersionedClause<OMPC_Device>,
     VersionedClause<OMPC_FirstPrivate>,
     VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_LastPrivate>,
     VersionedClause<OMPC_Map>,
     VersionedClause<OMPC_Private>,
     VersionedClause<OMPC_Reduction>,
     VersionedClause<OMPC_Shared>,
     VersionedClause<OMPC_UsesAllocators, 50>,
   ];
   let allowedOnceClauses = [