diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c
index b9024cb..eb4950c 100644
--- a/gcc/c-family/c-omp.c
+++ b/gcc/c-family/c-omp.c
@@ -1837,6 +1837,7 @@ c_omp_split_clauses (location_t loc, enum tree_code code,
 	case OMP_CLAUSE_DEVICE:
 	case OMP_CLAUSE_MAP:
 	case OMP_CLAUSE_IS_DEVICE_PTR:
+	case OMP_CLAUSE_HAS_DEVICE_ADDR:
 	case OMP_CLAUSE_DEFAULTMAP:
 	case OMP_CLAUSE_DEPEND:
 	  s = C_OMP_CLAUSE_SPLIT_TARGET;
diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index 0c5b07a..03baacd 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -89,8 +89,8 @@ enum pragma_kind {
 };
 
 
-/* All clauses defined by OpenACC 2.0, and OpenMP 2.5, 3.0, 3.1, 4.0, 4.5
-   and 5.0.  Used internally by both C and C++ parsers.  */
+/* All clauses defined by OpenACC 2.0, and OpenMP 2.5, 3.0, 3.1, 4.0, 4.5, 5.0,
+   and 5.1.  Used internally by both C and C++ parsers.  */
 enum pragma_omp_clause {
   PRAGMA_OMP_CLAUSE_NONE = 0,
 
@@ -114,6 +114,7 @@ enum pragma_omp_clause {
   PRAGMA_OMP_CLAUSE_FOR,
   PRAGMA_OMP_CLAUSE_FROM,
   PRAGMA_OMP_CLAUSE_GRAINSIZE,
+  PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR,
   PRAGMA_OMP_CLAUSE_HINT,
   PRAGMA_OMP_CLAUSE_IF,
   PRAGMA_OMP_CLAUSE_IN_REDUCTION,
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 869a811..e93a37f 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -12746,7 +12746,9 @@ c_parser_omp_clause_name (c_parser *parser)
 	    result = PRAGMA_OMP_CLAUSE_GRAINSIZE;
 	  break;
 	case 'h':
-	  if (!strcmp ("hint", p))
+	  if (!strcmp ("has_device_addr", p))
+	    result = PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR;
+	  else if (!strcmp ("hint", p))
 	    result = PRAGMA_OMP_CLAUSE_HINT;
 	  else if (!strcmp ("host", p))
 	    result = PRAGMA_OACC_CLAUSE_HOST;
@@ -14255,6 +14257,16 @@ c_parser_omp_clause_use_device_addr (c_parser *parser, tree list)
 				       list);
 }
 
+/* OpenMP 5.1:
+   has_device_addr ( variable-list ) */
+
+static tree
+c_parser_omp_clause_has_device_addr (c_parser *parser, tree list)
+{
+  return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_HAS_DEVICE_ADDR,
+				       list);
+}
+
 /* OpenMP 4.5:
    is_device_ptr ( variable-list ) */
 
@@ -16945,6 +16957,10 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask,
 	  clauses = c_parser_omp_clause_use_device_addr (parser, clauses);
 	  c_name = "use_device_addr";
 	  break;
+	case PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR:
+	  clauses = c_parser_omp_clause_has_device_addr (parser, clauses);
+	  c_name = "has_device_addr";
+	  break;
 	case PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR:
 	  clauses = c_parser_omp_clause_is_device_ptr (parser, clauses);
 	  c_name = "is_device_ptr";
@@ -20926,7 +20942,8 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser,
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_ALLOCATE)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEFAULTMAP)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IN_REDUCTION)	\
-	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR))
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR)\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR))
 
 static bool
 c_parser_omp_target (c_parser *parser, enum pragma_context context, bool *if_p)
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 0aac978..d677592 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -14054,7 +14054,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 {
   bitmap_head generic_head, firstprivate_head, lastprivate_head;
   bitmap_head aligned_head, map_head, map_field_head, map_firstprivate_head;
-  bitmap_head oacc_reduction_head;
+  bitmap_head oacc_reduction_head, has_device_addr_head, is_device_ptr_head;
   tree c, t, type, *pc;
   tree simdlen = NULL_TREE, safelen = NULL_TREE;
   bool branch_seen = false;
@@ -14089,6 +14089,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
   /* If ort == C_ORT_OMP used as nontemporal_head or use_device_xxx_head
      instead and for ort == C_ORT_OMP_TARGET used as in_reduction_head.  */
   bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack);
+  bitmap_initialize (&has_device_addr_head, &bitmap_default_obstack);
+  bitmap_initialize (&is_device_ptr_head, &bitmap_default_obstack);
 
   if (ort & C_ORT_ACC)
     for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
@@ -14517,7 +14519,9 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 			"%qE appears more than once in data clauses", t);
 	      remove = true;
 	    }
-	  else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE
+	  else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE
+		    || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR
+		    || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR)
 		   && bitmap_bit_p (&map_head, DECL_UID (t)))
 	    {
 	      if (ort == C_ORT_ACC)
@@ -15082,7 +15086,9 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 			"%qD appears more than once in data clauses", t);
 	      remove = true;
 	    }
-	  else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
+	  else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t))
+		   || bitmap_bit_p (&has_device_addr_head, DECL_UID (t))
+		   || bitmap_bit_p (&is_device_ptr_head, DECL_UID (t)))
 	    {
 	      if (ort == C_ORT_ACC)
 		error_at (OMP_CLAUSE_LOCATION (c),
@@ -15167,6 +15173,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	case OMP_CLAUSE_IS_DEVICE_PTR:
 	case OMP_CLAUSE_USE_DEVICE_PTR:
 	  t = OMP_CLAUSE_DECL (c);
+	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR)
+	    bitmap_set_bit (&is_device_ptr_head, DECL_UID (t));
 	  if (TREE_CODE (TREE_TYPE (t)) != POINTER_TYPE)
 	    {
 	      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR
@@ -15187,8 +15195,11 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	    }
 	  goto check_dup_generic;
 
+	case OMP_CLAUSE_HAS_DEVICE_ADDR:
 	case OMP_CLAUSE_USE_DEVICE_ADDR:
 	  t = OMP_CLAUSE_DECL (c);
+	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
+	    bitmap_set_bit (&has_device_addr_head, DECL_UID (t));
 	  if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
 	    c_mark_addressable (t);
 	  goto check_dup_generic;
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 0818d66..76582a7 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -36145,7 +36145,9 @@ cp_parser_omp_clause_name (cp_parser *parser)
 	    result = PRAGMA_OMP_CLAUSE_GRAINSIZE;
 	  break;
 	case 'h':
-	  if (!strcmp ("hint", p))
+	  if (!strcmp ("has_device_addr", p))
+	    result = PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR;
+	  else if (!strcmp ("hint", p))
 	    result = PRAGMA_OMP_CLAUSE_HINT;
 	  else if (!strcmp ("host", p))
 	    result = PRAGMA_OACC_CLAUSE_HOST;
@@ -39830,6 +39832,11 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask,
 					    clauses);
 	  c_name = "is_device_ptr";
 	  break;
+	case PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR:
+	  clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_HAS_DEVICE_ADDR,
+					    clauses);
+	  c_name = "has_device_addr";
+	  break;
 	case PRAGMA_OMP_CLAUSE_IF:
 	  clauses = cp_parser_omp_clause_if (parser, clauses, token->location,
 					     true);
@@ -44005,7 +44012,8 @@ cp_parser_omp_target_update (cp_parser *parser, cp_token *pragma_tok,
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEFAULTMAP)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_ALLOCATE)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IN_REDUCTION)	\
-	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR))
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR)\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR))
 
 static bool
 cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c
index 009fe6d..e5a6f2f 100644
--- a/gcc/cp/pt.c
+++ b/gcc/cp/pt.c
@@ -17442,6 +17442,7 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort,
 	case OMP_CLAUSE_USE_DEVICE_PTR:
 	case OMP_CLAUSE_USE_DEVICE_ADDR:
 	case OMP_CLAUSE_IS_DEVICE_PTR:
+	case OMP_CLAUSE_HAS_DEVICE_ADDR:
 	case OMP_CLAUSE_INCLUSIVE:
 	case OMP_CLAUSE_EXCLUSIVE:
 	  OMP_CLAUSE_DECL (nc)
@@ -17581,6 +17582,7 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort,
 	  case OMP_CLAUSE_USE_DEVICE_PTR:
 	  case OMP_CLAUSE_USE_DEVICE_ADDR:
 	  case OMP_CLAUSE_IS_DEVICE_PTR:
+	  case OMP_CLAUSE_HAS_DEVICE_ADDR:
 	  case OMP_CLAUSE_INCLUSIVE:
 	  case OMP_CLAUSE_EXCLUSIVE:
 	  case OMP_CLAUSE_ALLOCATE:
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 0d8e5fd..314d180 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -6580,7 +6580,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 {
   bitmap_head generic_head, firstprivate_head, lastprivate_head;
   bitmap_head aligned_head, map_head, map_field_head, map_firstprivate_head;
-  bitmap_head oacc_reduction_head;
+  bitmap_head oacc_reduction_head, has_device_addr_head, is_device_ptr_head;
   tree c, t, *pc;
   tree safelen = NULL_TREE;
   bool branch_seen = false;
@@ -6612,6 +6612,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
   /* If ort == C_ORT_OMP used as nontemporal_head or use_device_xxx_head
      instead and for ort == C_ORT_OMP_TARGET used as in_reduction_head.  */
   bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack);
+  bitmap_initialize (&has_device_addr_head, &bitmap_default_obstack);
+  bitmap_initialize (&is_device_ptr_head, &bitmap_default_obstack);
 
   if (ort & C_ORT_ACC)
     for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
@@ -6910,7 +6912,9 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 			"%qD appears more than once in data clauses", t);
 	      remove = true;
 	    }
-	  else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE
+	  else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE
+		    || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR
+		    || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR)
 		   && bitmap_bit_p (&map_head, DECL_UID (t)))
 	    {
 	      if (ort == C_ORT_ACC)
@@ -8032,7 +8036,9 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 			"%qD appears more than once in data clauses", t);
 	      remove = true;
 	    }
-	  else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
+	  else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t))
+		   || bitmap_bit_p (&has_device_addr_head, DECL_UID (t))
+		   || bitmap_bit_p (&is_device_ptr_head, DECL_UID (t)))
 	    {
 	      if (ort == C_ORT_ACC)
 		error_at (OMP_CLAUSE_LOCATION (c),
@@ -8286,6 +8292,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	case OMP_CLAUSE_USE_DEVICE_PTR:
 	  field_ok = (ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP;
 	  t = OMP_CLAUSE_DECL (c);
+	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR)
+	    bitmap_set_bit (&is_device_ptr_head, DECL_UID (t));
 	  if (!type_dependent_expression_p (t))
 	    {
 	      tree type = TREE_TYPE (t);
@@ -8315,9 +8323,12 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	    }
 	  goto check_dup_generic;
 
+	case OMP_CLAUSE_HAS_DEVICE_ADDR:
 	case OMP_CLAUSE_USE_DEVICE_ADDR:
 	  field_ok = true;
 	  t = OMP_CLAUSE_DECL (c);
+	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
+	    bitmap_set_bit (&has_device_addr_head, DECL_UID (t));
 	  if (!processing_template_decl
 	      && (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
 	      && !TYPE_REF_P (TREE_TYPE (t))
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index d8e4b13..3c0b262 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -10018,6 +10018,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	  flags = GOVD_EXPLICIT;
 	  goto do_add;
 
+	case OMP_CLAUSE_HAS_DEVICE_ADDR:
 	case OMP_CLAUSE_IS_DEVICE_PTR:
 	  flags = GOVD_FIRSTPRIVATE | GOVD_EXPLICIT;
 	  goto do_add;
@@ -11446,6 +11447,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 	case OMP_CLAUSE_DETACH:
 	case OMP_CLAUSE_USE_DEVICE_PTR:
 	case OMP_CLAUSE_USE_DEVICE_ADDR:
+	case OMP_CLAUSE_HAS_DEVICE_ADDR:
 	case OMP_CLAUSE_IS_DEVICE_PTR:
 	case OMP_CLAUSE_ASYNC:
 	case OMP_CLAUSE_WAIT:
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 057b7ae..3c92cb2 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -1368,7 +1368,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	  decl = OMP_CLAUSE_DECL (c);
 	do_private:
 	  if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE
-	       || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR)
+	       || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR
+	       || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
 	      && is_gimple_omp_offloaded (ctx->stmt))
 	    {
 	      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
@@ -1376,7 +1377,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 		  by_ref = !omp_privatize_by_reference (decl);
 		  install_var_field (decl, by_ref, 3, ctx);
 		}
-	      else if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
+	      else if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE
+		       || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
 		install_var_field (decl, true, 3, ctx);
 	      else
 		install_var_field (decl, false, 3, ctx);
@@ -1446,6 +1448,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	  break;
 
 	case OMP_CLAUSE_IS_DEVICE_PTR:
+	case OMP_CLAUSE_HAS_DEVICE_ADDR:
 	  decl = OMP_CLAUSE_DECL (c);
 	  goto do_private;
 
@@ -1722,12 +1725,14 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	case OMP_CLAUSE_FIRSTPRIVATE:
 	case OMP_CLAUSE_PRIVATE:
 	case OMP_CLAUSE_LINEAR:
+	case OMP_CLAUSE_HAS_DEVICE_ADDR:
 	case OMP_CLAUSE_IS_DEVICE_PTR:
 	  decl = OMP_CLAUSE_DECL (c);
 	  if (is_variable_sized (decl))
 	    {
 	      if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE
-		   || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR)
+		   || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR
+		   || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
 		  && is_gimple_omp_offloaded (ctx->stmt))
 		{
 		  tree decl2 = DECL_VALUE_EXPR (decl);
@@ -12791,6 +12796,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 
       case OMP_CLAUSE_USE_DEVICE_PTR:
       case OMP_CLAUSE_USE_DEVICE_ADDR:
+      case OMP_CLAUSE_HAS_DEVICE_ADDR:
       case OMP_CLAUSE_IS_DEVICE_PTR:
 	var = OMP_CLAUSE_DECL (c);
 	map_cnt++;
@@ -12807,7 +12813,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	    SET_DECL_VALUE_EXPR (new_var, x);
 	    DECL_HAS_VALUE_EXPR_P (new_var) = 1;
 	  }
-	else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR
+	else if (((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR
+		   || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
 		  && !omp_privatize_by_reference (var)
 		  && !omp_is_allocatable_or_ptr (var)
 		  && !lang_hooks.decls.omp_array_data (var, true))
@@ -13260,17 +13267,20 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 
 	  case OMP_CLAUSE_USE_DEVICE_PTR:
 	  case OMP_CLAUSE_USE_DEVICE_ADDR:
+	  case OMP_CLAUSE_HAS_DEVICE_ADDR:
 	  case OMP_CLAUSE_IS_DEVICE_PTR:
 	    ovar = OMP_CLAUSE_DECL (c);
 	    var = lookup_decl_in_outer_ctx (ovar, ctx);
 
 	    if (lang_hooks.decls.omp_array_data (ovar, true))
 	      {
-		tkind = (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR
+		tkind = ((OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR
+			  && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_HAS_DEVICE_ADDR)
 			 ? GOMP_MAP_USE_DEVICE_PTR : GOMP_MAP_FIRSTPRIVATE_INT);
 		x = build_sender_ref ((splay_tree_key) &DECL_NAME (ovar), ctx);
 	      }
-	    else if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR)
+	    else if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR
+		     && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_HAS_DEVICE_ADDR)
 	      {
 		tkind = GOMP_MAP_USE_DEVICE_PTR;
 		x = build_sender_ref ((splay_tree_key) &DECL_UID (ovar), ctx);
@@ -13292,7 +13302,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	    type = TREE_TYPE (ovar);
 	    if (lang_hooks.decls.omp_array_data (ovar, true))
 	      var = lang_hooks.decls.omp_array_data (ovar, false);
-	    else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR
+	    else if (((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR
+		      || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
 		      && !omp_privatize_by_reference (ovar)
 		      && !omp_is_allocatable_or_ptr (ovar))
 		     || TREE_CODE (type) == ARRAY_TYPE)
@@ -13307,6 +13318,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 		    if (POINTER_TYPE_P (type)
 			&& TREE_CODE (type) != ARRAY_TYPE
 			&& ((OMP_CLAUSE_CODE (c) != OMP_CLAUSE_USE_DEVICE_ADDR
+			    && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_HAS_DEVICE_ADDR
 			    && !omp_is_allocatable_or_ptr (ovar))
 			   || (omp_privatize_by_reference (ovar)
 			       && omp_is_allocatable_or_ptr (ovar))))
@@ -13504,6 +13516,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	    break;
 	  case OMP_CLAUSE_USE_DEVICE_PTR:
 	  case OMP_CLAUSE_USE_DEVICE_ADDR:
+	  case OMP_CLAUSE_HAS_DEVICE_ADDR:
 	  case OMP_CLAUSE_IS_DEVICE_PTR:
 	    tree new_var;
 	    gimple_seq assign_body;
@@ -13514,7 +13527,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	    var = OMP_CLAUSE_DECL (c);
 	    is_array_data = lang_hooks.decls.omp_array_data (var, true) != NULL;
 
-	    if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR)
+	    if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR
+		&& OMP_CLAUSE_CODE (c) != OMP_CLAUSE_HAS_DEVICE_ADDR)
 	      x = build_sender_ref (is_array_data
 				    ? (splay_tree_key) &DECL_NAME (var)
 				    : (splay_tree_key) &DECL_UID (var), ctx);
@@ -13566,7 +13580,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 		gimple_seq_add_stmt (&assign_body,
 				     gimple_build_assign (new_var, x));
 	      }
-	    else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR
+	    else if (((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR
+		      || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
 		      && !omp_privatize_by_reference (var)
 		      && !omp_is_allocatable_or_ptr (var))
 		     || TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
diff --git a/gcc/testsuite/c-c++-common/gomp/clauses-1.c b/gcc/testsuite/c-c++-common/gomp/clauses-1.c
index 742132f..9da5255 100644
--- a/gcc/testsuite/c-c++-common/gomp/clauses-1.c
+++ b/gcc/testsuite/c-c++-common/gomp/clauses-1.c
@@ -102,7 +102,7 @@ baz (int d, int m, int i1, int i2, int p, int *idp, int s,
 }
 
 void
-bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
+bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int hda, int s,
      int nte, int tl, int nth, int g, int nta, int fi, int pp, int *q, int *dd, int ntm)
 {
   #pragma omp for simd \
@@ -138,20 +138,20 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
   #pragma omp target parallel \
     device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \
     if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) \
-    nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2)
+    nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr(hda)
     ;
   #pragma omp target parallel for \
     device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \
     if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) \
     lastprivate (l) linear (ll:1) ordered schedule(static, 4) collapse(1) nowait depend(inout: dd[0]) \
-    allocate (omp_default_mem_alloc:f) in_reduction(+:r2)
+    allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr(hda)
   for (int i = 0; i < 64; i++)
     ll++;
   #pragma omp target parallel for \
     device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \
     if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) \
     lastprivate (l) linear (ll:1) schedule(static, 4) collapse(1) nowait depend(inout: dd[0]) order(concurrent) \
-    allocate (omp_default_mem_alloc:f) in_reduction(+:r2)
+    allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr(hda)
   for (int i = 0; i < 64; i++)
     ll++;
   #pragma omp target parallel for simd \
@@ -159,18 +159,19 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
     if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) \
     lastprivate (l) linear (ll:1) schedule(static, 4) collapse(1) \
     safelen(8) simdlen(4) aligned(q: 32) nowait depend(inout: dd[0]) nontemporal(ntm) if (simd: i3) order(concurrent) \
-    allocate (omp_default_mem_alloc:f) in_reduction(+:r2)
+    allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr(hda)
   for (int i = 0; i < 64; i++)
     ll++;
   #pragma omp target teams \
     device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \
     shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) nowait depend(inout: dd[0]) \
-    allocate (omp_default_mem_alloc:f) in_reduction(+:r2)
+    allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr(hda)
     ;
   #pragma omp target teams distribute \
     device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \
     shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) order(concurrent) \
-    collapse(1) dist_schedule(static, 16) nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2)
+    collapse(1) dist_schedule(static, 16) nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2) \
+    has_device_addr(hda)
   for (int i = 0; i < 64; i++)
     ;
   #pragma omp target teams distribute parallel for \
@@ -179,7 +180,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
     collapse(1) dist_schedule(static, 16) \
     if (parallel: i2) num_threads (nth) proc_bind(spread) \
     lastprivate (l) schedule(static, 4) nowait depend(inout: dd[0]) order(concurrent) \
-     allocate (omp_default_mem_alloc:f) in_reduction(+:r2)
+     allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr(hda)
   for (int i = 0; i < 64; i++)
     ll++;
   #pragma omp target teams distribute parallel for simd \
@@ -189,7 +190,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
     if (parallel: i2) num_threads (nth) proc_bind(spread) \
     lastprivate (l) schedule(static, 4) order(concurrent) \
     safelen(8) simdlen(4) aligned(q: 32) nowait depend(inout: dd[0]) nontemporal(ntm) if (simd: i3) \
-    allocate (omp_default_mem_alloc:f) in_reduction(+:r2)
+    allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr(hda)
   for (int i = 0; i < 64; i++)
     ll++;
   #pragma omp target teams distribute simd \
@@ -197,14 +198,14 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
     shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) \
     collapse(1) dist_schedule(static, 16) order(concurrent) \
     safelen(8) simdlen(4) aligned(q: 32) nowait depend(inout: dd[0]) nontemporal(ntm) \
-    allocate (omp_default_mem_alloc:f) in_reduction(+:r2)
+    allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr(hda)
   for (int i = 0; i < 64; i++)
     ll++;
   #pragma omp target simd \
     device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \
     safelen(8) simdlen(4) lastprivate (l) linear(ll: 1) aligned(q: 32) reduction(+:r) \
     nowait depend(inout: dd[0]) nontemporal(ntm) if(simd:i3) order(concurrent) \
-    allocate (omp_default_mem_alloc:f) in_reduction(+:r2)
+    allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr(hda)
   for (int i = 0; i < 64; i++)
     ll++;
   #pragma omp taskgroup task_reduction(+:r2) allocate (r2)
@@ -430,28 +431,28 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
     device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \
     if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) \
     nowait depend(inout: dd[0]) lastprivate (l) bind(parallel) order(concurrent) collapse(1) \
-    allocate (omp_default_mem_alloc: f) in_reduction(+:r2)
+    allocate (omp_default_mem_alloc: f) in_reduction(+:r2) has_device_addr(hda)
   for (l = 0; l < 64; ++l)
     ;
   #pragma omp target parallel loop \
     device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \
     if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) \
     nowait depend(inout: dd[0]) lastprivate (l) order(concurrent) collapse(1) \
-    allocate (omp_default_mem_alloc: f) in_reduction(+:r2)
+    allocate (omp_default_mem_alloc: f) in_reduction(+:r2) has_device_addr(hda)
   for (l = 0; l < 64; ++l)
     ;
   #pragma omp target teams loop \
     device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \
     shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) nowait depend(inout: dd[0]) \
     lastprivate (l) bind(teams) collapse(1) \
-    allocate (omp_default_mem_alloc: f) in_reduction(+:r2)
+    allocate (omp_default_mem_alloc: f) in_reduction(+:r2) has_device_addr(hda)
   for (l = 0; l < 64; ++l)
     ;
   #pragma omp target teams loop \
     device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \
     shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) nowait depend(inout: dd[0]) \
     lastprivate (l) order(concurrent) collapse(1) \
-    allocate (omp_default_mem_alloc: f) in_reduction(+:r2)
+    allocate (omp_default_mem_alloc: f) in_reduction(+:r2) has_device_addr(hda)
   for (l = 0; l < 64; ++l)
     ;
 }
diff --git a/gcc/testsuite/c-c++-common/gomp/target-has-device-addr-1.c b/gcc/testsuite/c-c++-common/gomp/target-has-device-addr-1.c
new file mode 100644
index 0000000..3a12e62
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-has-device-addr-1.c
@@ -0,0 +1,27 @@
+/* { dg-do compile } */
+
+void
+foo ()
+{
+  int * x;
+  #pragma omp target is_device_ptr(x) has_device_addr(x) /*{ dg-error "'x' appears more than once in data clauses" } */
+  ;
+  #pragma omp target has_device_addr(x) is_device_ptr(x) /* { dg-error "'x' appears more than once in data clauses" } */
+  ;
+
+  int y = 42;
+  #pragma omp target has_device_addr(y) has_device_addr(y) /* { dg-error "'y' appears more than once in data clauses" } */
+  ;
+
+  #pragma omp target private(y) has_device_addr(y) /*{ dg-error "'y' appears more than once in data clauses" } */
+  ;
+  #pragma omp target has_device_addr(y) private(y) /*{ dg-error "'y' appears more than once in data clauses" } */
+  ;
+  #pragma omp target firstprivate(y) has_device_addr(y) /*{ dg-error "'y' appears more than once in data clauses" } */
+  ;
+
+  #pragma omp target has_device_addr(y) map(y) /* { dg-error "'y' appears both in data and map clauses" } */
+  ;
+  #pragma omp target map(y) has_device_addr(y) /* { dg-error "'y' appears both in data and map clauses" } */
+  ;
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/target-is-device-ptr.c b/gcc/testsuite/c-c++-common/gomp/target-is-device-ptr.c
new file mode 100644
index 0000000..ecf30ca
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-is-device-ptr.c
@@ -0,0 +1,22 @@
+/* { dg-do compile } */
+
+void
+foo ()
+{
+  int *x;
+
+  #pragma omp target is_device_ptr(x) is_device_ptr(x) /* { dg-error "'x' appears more than once in data clauses" } */
+  ;
+
+  #pragma omp target private(x) is_device_ptr(x) /*{ dg-error "'x' appears more than once in data clauses" } */
+  ;
+  #pragma omp target is_device_ptr(x) private(x) /*{ dg-error "'x' appears more than once in data clauses" } */
+  ;
+  #pragma omp target firstprivate(x) is_device_ptr(x) /*{ dg-error "'x' appears more than once in data clauses" } */
+  ;
+
+  #pragma omp target is_device_ptr(x) map(x) /* { dg-error "'x' appears both in data and map clauses" } */
+  ;
+  #pragma omp target map(x) is_device_ptr(x) /* { dg-error "'x' appears both in data and map clauses" } */
+  ;
+}
diff --git a/gcc/testsuite/g++.dg/gomp/attrs-1.C b/gcc/testsuite/g++.dg/gomp/attrs-1.C
index 2a5f2cf..95a8fc6a 100644
--- a/gcc/testsuite/g++.dg/gomp/attrs-1.C
+++ b/gcc/testsuite/g++.dg/gomp/attrs-1.C
@@ -121,7 +121,7 @@ baz (int d, int m, int i1, int i2, int p, int *idp, int s,
 }
 
 void
-bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
+bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int hda, int s,
      int nte, int tl, int nth, int g, int nta, int fi, int pp, int *q, int *dd, int ntm,
      const char *msg)
 {
@@ -185,20 +185,20 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
   [[omp::directive (target parallel
     device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp)
     if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread)
-    nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2))]]
+    nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr (hda))]]
     ;
   [[omp::directive (target parallel for
     device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp)
     if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread)
     lastprivate (l) linear (ll:1) ordered schedule(static, 4) collapse(1) nowait depend(inout: dd[0])
-    allocate (omp_default_mem_alloc:f) in_reduction(+:r2))]]
+    allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr (hda))]]
   for (int i = 0; i < 64; i++)
     ll++;
   [[omp::directive (target parallel for
     device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp)
     if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread)
     lastprivate (l) linear (ll:1) schedule(static, 4) collapse(1) nowait depend(inout: dd[0]) order(concurrent)
-    allocate (omp_default_mem_alloc:f) in_reduction(+:r2))]]
+    allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr (hda))]]
   for (int i = 0; i < 64; i++)
     ll++;
   [[omp::sequence (omp::directive (target parallel for simd
@@ -206,22 +206,23 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
     if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread)
     lastprivate (l) linear (ll:1) schedule(static, 4) collapse(1)
     safelen(8) simdlen(4) aligned(q: 32) nowait depend(inout: dd[0]) nontemporal(ntm) if (simd: i3) order(concurrent)
-    allocate (omp_default_mem_alloc:f) in_reduction(+:r2)))]]
+    allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr (hda)))]]
   for (int i = 0; i < 64; i++)
     ll++;
   [[omp::sequence (directive (target teams
     device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp)
     shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) nowait depend(inout: dd[0])
-    allocate (omp_default_mem_alloc:f) in_reduction(+:r2)))]]
+    allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr (hda)))]]
     ;
   [[omp::sequence (directive (target
     device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp)
-    nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2)))]]
+    nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr (hda)))]]
     ;
   [[omp::sequence (omp::directive (target teams distribute
     device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp)
     shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) order(concurrent)
-    collapse(1) dist_schedule(static, 16) nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2)))]]
+    collapse(1) dist_schedule(static, 16) nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2)
+    has_device_addr (hda)))]]
   for (int i = 0; i < 64; i++)
     ;
   [[omp::directive (target teams distribute parallel for
@@ -230,7 +231,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
     collapse(1) dist_schedule(static, 16)
     if (parallel: i2) num_threads (nth) proc_bind(spread)
     lastprivate (l) schedule(static, 4) nowait depend(inout: dd[0]) order(concurrent)
-     allocate (omp_default_mem_alloc:f) in_reduction(+:r2))]]
+     allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr (hda))]]
   for (int i = 0; i < 64; i++)
     ll++;
   [[omp::directive (target teams distribute parallel for simd
@@ -240,7 +241,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
     if (parallel: i2) num_threads (nth) proc_bind(spread)
     lastprivate (l) schedule(static, 4) order(concurrent)
     safelen(8) simdlen(4) aligned(q: 32) nowait depend(inout: dd[0]) nontemporal(ntm) if (simd: i3)
-    allocate (omp_default_mem_alloc:f) in_reduction(+:r2))]]
+    allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr (hda))]]
   for (int i = 0; i < 64; i++)
     ll++;
   [[omp::directive (target teams distribute simd
@@ -248,14 +249,14 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
     shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl)
     collapse(1) dist_schedule(static, 16) order(concurrent)
     safelen(8) simdlen(4) aligned(q: 32) nowait depend(inout: dd[0]) nontemporal(ntm)
-    allocate (omp_default_mem_alloc:f) in_reduction(+:r2))]]
+    allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr (hda))]]
   for (int i = 0; i < 64; i++)
     ll++;
   [[omp::directive (target simd
     device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp)
     safelen(8) simdlen(4) lastprivate (l) linear(ll: 1) aligned(q: 32) reduction(+:r)
     nowait depend(inout: dd[0]) nontemporal(ntm) if(simd:i3) order(concurrent)
-    allocate (omp_default_mem_alloc:f) in_reduction(+:r2))]]
+    allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr (hda))]]
   for (int i = 0; i < 64; i++)
     ll++;
   [[omp::sequence (directive (taskgroup task_reduction(+:r2) allocate (r2)),
@@ -515,28 +516,28 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
     device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp)
     if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread)
     nowait depend(inout: dd[0]) lastprivate (l) bind(parallel) order(concurrent) collapse(1)
-    allocate (omp_default_mem_alloc: f) in_reduction(+:r2))]]
+    allocate (omp_default_mem_alloc: f) in_reduction(+:r2) has_device_addr (hda))]]
   for (l = 0; l < 64; ++l)
     ;
   [[omp::directive (target parallel loop
     device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp)
     if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread)
     nowait depend(inout: dd[0]) lastprivate (l) order(concurrent) collapse(1)
-    allocate (omp_default_mem_alloc: f) in_reduction(+:r2))]]
+    allocate (omp_default_mem_alloc: f) in_reduction(+:r2) has_device_addr (hda))]]
   for (l = 0; l < 64; ++l)
     ;
   [[omp::directive (target teams loop
     device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp)
     shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) nowait depend(inout: dd[0])
     lastprivate (l) bind(teams) collapse(1)
-    allocate (omp_default_mem_alloc: f) in_reduction(+:r2))]]
+    allocate (omp_default_mem_alloc: f) in_reduction(+:r2) has_device_addr (hda))]]
   for (l = 0; l < 64; ++l)
     ;
   [[omp::directive (target teams loop
     device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp)
     shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) nowait depend(inout: dd[0])
     lastprivate (l) order(concurrent) collapse(1)
-    allocate (omp_default_mem_alloc: f) in_reduction(+:r2))]]
+    allocate (omp_default_mem_alloc: f) in_reduction(+:r2) has_device_addr (hda))]]
   for (l = 0; l < 64; ++l)
     ;
   [[omp::directive (critical)]] {
diff --git a/gcc/testsuite/g++.dg/gomp/attrs-2.C b/gcc/testsuite/g++.dg/gomp/attrs-2.C
index c00be7f..36a7584 100644
--- a/gcc/testsuite/g++.dg/gomp/attrs-2.C
+++ b/gcc/testsuite/g++.dg/gomp/attrs-2.C
@@ -121,7 +121,7 @@ baz (int d, int m, int i1, int i2, int p, int *idp, int s,
 }
 
 void
-bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
+bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int hda, int s,
      int nte, int tl, int nth, int g, int nta, int fi, int pp, int *q, int *dd, int ntm,
      const char *msg)
 {
@@ -185,20 +185,20 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
   [[omp::directive (target parallel,
     device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp),
     if (parallel: i2),default(shared),shared(s),reduction(+:r),num_threads (nth),proc_bind(spread)
-    nowait depend(inout: dd[0]),allocate (omp_default_mem_alloc:f),in_reduction(+:r2))]]
+    nowait depend(inout: dd[0]),allocate (omp_default_mem_alloc:f),in_reduction(+:r2),has_device_addr (hda))]]
     ;
   [[omp::directive (target parallel for,
     device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp),
     if (parallel: i2),default(shared),shared(s),reduction(+:r),num_threads (nth),proc_bind(spread),
     lastprivate (l),linear (ll:1),ordered schedule(static, 4),collapse(1),nowait depend(inout: dd[0]),
-    allocate (omp_default_mem_alloc:f),in_reduction(+:r2))]]
+    allocate (omp_default_mem_alloc:f),in_reduction(+:r2),has_device_addr (hda))]]
   for (int i = 0; i < 64; i++)
     ll++;
   [[using omp:directive (target parallel for,
     device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp),
     if (parallel: i2),default(shared),shared(s),reduction(+:r),num_threads (nth),proc_bind(spread),
     lastprivate (l),linear (ll:1),schedule(static, 4),collapse(1),nowait depend(inout: dd[0]),order(concurrent),
-    allocate (omp_default_mem_alloc:f),in_reduction(+:r2))]]
+    allocate (omp_default_mem_alloc:f),in_reduction(+:r2),has_device_addr (hda))]]
   for (int i = 0; i < 64; i++)
     ll++;
   [[omp::sequence (omp::directive (target parallel for simd,
@@ -206,22 +206,23 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
     if (parallel: i2),default(shared),shared(s),reduction(+:r),num_threads (nth),proc_bind(spread),
     lastprivate (l),linear (ll:1),schedule(static, 4),collapse(1),
     safelen(8),simdlen(4),aligned(q: 32),nowait depend(inout: dd[0]),nontemporal(ntm),if (simd: i3),order(concurrent),
-    allocate (omp_default_mem_alloc:f),in_reduction(+:r2)))]]
+    allocate (omp_default_mem_alloc:f),in_reduction(+:r2),has_device_addr (hda)))]]
   for (int i = 0; i < 64; i++)
     ll++;
   [[using omp:sequence (directive (target teams,
     device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp),
-    shared(s),default(shared),reduction(+:r),num_teams(nte),thread_limit(tl),nowait, depend(inout: dd[0]),
-    allocate (omp_default_mem_alloc:f) in_reduction(+:r2)))]]
+    shared(s),default(shared),reduction(+:r),num_teams(nte),thread_limit(tl),nowait,depend(inout: dd[0]),
+    allocate (omp_default_mem_alloc:f),in_reduction(+:r2),has_device_addr (hda)))]]
     ;
   [[using omp:sequence (directive (target,
     device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp),
-    nowait depend(inout: dd[0]),allocate (omp_default_mem_alloc:f),in_reduction(+:r2)))]]
+    nowait depend(inout: dd[0]),allocate (omp_default_mem_alloc:f),in_reduction(+:r2),has_device_addr(hda)))]]
     ;
   [[omp::sequence (omp::directive (target teams distribute,
     device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp),
     shared(s),default(shared),reduction(+:r),num_teams(nte),thread_limit(tl),order(concurrent),
-    collapse(1),dist_schedule(static, 16),nowait depend(inout: dd[0]),allocate (omp_default_mem_alloc:f),in_reduction(+:r2)))]]
+    collapse(1),dist_schedule(static, 16),nowait depend(inout: dd[0]),allocate (omp_default_mem_alloc:f),in_reduction(+:r2),
+    has_device_addr (hda)))]]
   for (int i = 0; i < 64; i++)
     ;
   [[omp::directive (target teams distribute parallel for,
@@ -230,7 +231,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
     collapse(1),dist_schedule(static, 16),
     if (parallel: i2),num_threads (nth),proc_bind(spread),
     lastprivate (l),schedule(static, 4),nowait depend(inout: dd[0]),order(concurrent),
-     allocate (omp_default_mem_alloc:f),in_reduction(+:r2))]]
+    allocate (omp_default_mem_alloc:f),in_reduction(+:r2),has_device_addr (hda))]]
   for (int i = 0; i < 64; i++)
     ll++;
   [[omp::directive (target teams distribute parallel for simd,
@@ -240,7 +241,7 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
     if (parallel: i2),num_threads (nth),proc_bind(spread),
     lastprivate (l),schedule(static, 4),order(concurrent),
     safelen(8),simdlen(4),aligned(q: 32),nowait depend(inout: dd[0]),nontemporal(ntm),if (simd: i3),
-    allocate (omp_default_mem_alloc:f),in_reduction(+:r2))]]
+    allocate (omp_default_mem_alloc:f),in_reduction(+:r2),has_device_addr (hda))]]
   for (int i = 0; i < 64; i++)
     ll++;
   [[omp::directive (target teams distribute simd,
@@ -248,14 +249,14 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
     shared(s),default(shared),reduction(+:r),num_teams(nte),thread_limit(tl),
     collapse(1),dist_schedule(static, 16),order(concurrent),
     safelen(8),simdlen(4),aligned(q: 32),nowait depend(inout: dd[0]),nontemporal(ntm),
-    allocate (omp_default_mem_alloc:f),in_reduction(+:r2))]]
+    allocate (omp_default_mem_alloc:f),in_reduction(+:r2),has_device_addr (hda))]]
   for (int i = 0; i < 64; i++)
     ll++;
   [[omp::directive (target simd,
     device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp),
     safelen(8),simdlen(4),lastprivate (l),linear(ll: 1),aligned(q: 32),reduction(+:r),
     nowait depend(inout: dd[0]),nontemporal(ntm),if(simd:i3),order(concurrent),
-    allocate (omp_default_mem_alloc:f),in_reduction(+:r2))]]
+    allocate (omp_default_mem_alloc:f),in_reduction(+:r2),has_device_addr (hda))]]
   for (int i = 0; i < 64; i++)
     ll++;
   [[omp::sequence (directive (taskgroup, task_reduction(+:r2), allocate (r2)),
@@ -515,28 +516,28 @@ bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
     device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp),
     if (parallel: i2),default(shared),shared(s),reduction(+:r),num_threads (nth),proc_bind(spread),
     nowait depend(inout: dd[0]),lastprivate (l),bind(parallel),order(concurrent),collapse(1),
-    allocate (omp_default_mem_alloc: f),in_reduction(+:r2))]]
+    allocate (omp_default_mem_alloc: f),in_reduction(+:r2),has_device_addr (hda))]]
   for (l = 0; l < 64; ++l)
     ;
   [[omp::directive (target parallel loop,
     device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp),
     if (parallel: i2),default(shared),shared(s),reduction(+:r),num_threads (nth),proc_bind(spread),
     nowait depend(inout: dd[0]),lastprivate (l),order(concurrent),collapse(1),
-    allocate (omp_default_mem_alloc: f),in_reduction(+:r2))]]
+    allocate (omp_default_mem_alloc: f),in_reduction(+:r2),has_device_addr (hda))]]
   for (l = 0; l < 64; ++l)
     ;
   [[omp::directive (target teams loop,
     device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp),
     shared(s),default(shared),reduction(+:r),num_teams(nte),thread_limit(tl),nowait,depend(inout: dd[0]),
     lastprivate (l),bind(teams),collapse(1),
-    allocate (omp_default_mem_alloc: f),in_reduction(+:r2))]]
+    allocate (omp_default_mem_alloc: f),in_reduction(+:r2),has_device_addr (hda))]]
   for (l = 0; l < 64; ++l)
     ;
   [[omp::directive (target teams loop,
     device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp),
     shared(s),default(shared),reduction(+:r),num_teams(nte),thread_limit(tl),nowait,depend(inout: dd[0]),
     lastprivate (l),order(concurrent),collapse(1)
-    allocate (omp_default_mem_alloc: f),in_reduction(+:r2))]]
+    allocate (omp_default_mem_alloc: f),in_reduction(+:r2),has_device_addr (hda))]]
   for (l = 0; l < 64; ++l)
     ;
   [[omp::directive (critical)]] {
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index d3d2a8d..9091b9e 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -321,6 +321,9 @@ enum omp_clause_code {
   /* OpenMP clause: is_device_ptr (variable-list).  */
   OMP_CLAUSE_IS_DEVICE_PTR,
 
+  /* OpenMP clause: has_device_addr (variable-list).  */
+  OMP_CLAUSE_HAS_DEVICE_ADDR,
+
   /* OpenMP clause: inclusive (variable-list).  */
   OMP_CLAUSE_INCLUSIVE,
 
diff --git a/gcc/tree-nested.c b/gcc/tree-nested.c
index c7f50eb..449e4be 100644
--- a/gcc/tree-nested.c
+++ b/gcc/tree-nested.c
@@ -1339,6 +1339,7 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
 	case OMP_CLAUSE_LINK:
 	case OMP_CLAUSE_USE_DEVICE_PTR:
 	case OMP_CLAUSE_USE_DEVICE_ADDR:
+	case OMP_CLAUSE_HAS_DEVICE_ADDR:
 	case OMP_CLAUSE_IS_DEVICE_PTR:
 	case OMP_CLAUSE_DETACH:
 	do_decl_clause:
@@ -2123,6 +2124,7 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
 	case OMP_CLAUSE_LINK:
 	case OMP_CLAUSE_USE_DEVICE_PTR:
 	case OMP_CLAUSE_USE_DEVICE_ADDR:
+	case OMP_CLAUSE_HAS_DEVICE_ADDR:
 	case OMP_CLAUSE_IS_DEVICE_PTR:
 	case OMP_CLAUSE_DETACH:
 	do_decl_clause:
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 275dc7d..676fadd 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -493,6 +493,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
     case OMP_CLAUSE_USE_DEVICE_ADDR:
       name = "use_device_addr";
       goto print_remap;
+    case OMP_CLAUSE_HAS_DEVICE_ADDR:
+      name = "has_device_addr";
+      goto print_remap;
     case OMP_CLAUSE_IS_DEVICE_PTR:
       name = "is_device_ptr";
       goto print_remap;
diff --git a/gcc/tree.c b/gcc/tree.c
index 7bfd641..a47375e 100644
--- a/gcc/tree.c
+++ b/gcc/tree.c
@@ -301,6 +301,7 @@ unsigned const char omp_clause_num_ops[] =
   1, /* OMP_CLAUSE_USE_DEVICE_PTR  */
   1, /* OMP_CLAUSE_USE_DEVICE_ADDR  */
   1, /* OMP_CLAUSE_IS_DEVICE_PTR  */
+  1, /* OMP_CLAUSE_HAS_DEVICE_ADDR  */
   1, /* OMP_CLAUSE_INCLUSIVE  */
   1, /* OMP_CLAUSE_EXCLUSIVE  */
   2, /* OMP_CLAUSE_FROM  */
@@ -390,6 +391,7 @@ const char * const omp_clause_code_name[] =
   "use_device_ptr",
   "use_device_addr",
   "is_device_ptr",
+  "has_device_addr",
   "inclusive",
   "exclusive",
   "from",
diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
index bdd7e3a..44d060c 100644
--- a/libgomp/libgomp.texi
+++ b/libgomp/libgomp.texi
@@ -293,7 +293,8 @@ The OpenMP 4.5 specification is fully supported.
 @item @code{align} clause/modifier in @code{allocate} directive/clause
       and @code{allocator} directive @tab P @tab C/C++ on clause only
 @item @code{thread_limit} clause to @code{target} construct @tab N @tab
-@item @code{has_device_addr} clause to @code{target} construct @tab N @tab
+@item @code{has_device_addr} clause to @code{target} construct @tab P
+      @tab C/C++ on clause only
 @item iterators in @code{target update} motion clauses and @code{map}
       clauses @tab N @tab
 @item indirect calls to the device version of a procedure or function in
diff --git a/libgomp/testsuite/libgomp.c++/target-has-device-addr-2.C b/libgomp/testsuite/libgomp.c++/target-has-device-addr-2.C
new file mode 100644
index 0000000..b93b880
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-has-device-addr-2.C
@@ -0,0 +1,24 @@
+/* Testing the 'has_device_addr' clause on the target construct without
+   enclosing 'target data' construct. */
+
+#include <omp.h>
+
+int
+main ()
+{
+  int *dp = (int*)omp_target_alloc(sizeof(int), 0);
+
+  #pragma omp target is_device_ptr(dp)
+    *dp = 42;
+
+  int &x = *dp;
+
+  #pragma omp target has_device_addr(x)
+    x = 24;
+
+  #pragma omp target has_device_addr(x)
+    if (x != 24)
+      __builtin_abort ();
+
+  omp_target_free(dp, 0);
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-has-device-addr-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-has-device-addr-1.c
new file mode 100644
index 0000000..76739d4
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-has-device-addr-1.c
@@ -0,0 +1,23 @@
+/* Testing the 'has_device_addr' clause on the target construct with
+   enclosing 'target data' construct. */
+
+int
+main ()
+{
+  int x = 24;
+
+  #pragma omp target data map(x) use_device_addr(x)
+    #pragma omp target has_device_addr(x)
+      x = 42;
+  if (x != 42)
+    __builtin_abort ();
+
+  x = 24;
+  #pragma omp target data map(to: x) use_device_addr(x)
+    #pragma omp target has_device_addr(x)
+      x = 42;
+  if (x != 24)
+    __builtin_abort ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-has-device-addr-3.c b/libgomp/testsuite/libgomp.c-c++-common/target-has-device-addr-3.c
new file mode 100644
index 0000000..d90f389
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-has-device-addr-3.c
@@ -0,0 +1,30 @@
+/* Testing the 'has_device_addr' clause on the target construct with
+   enclosing 'target data' construct. */
+
+#include <stdlib.h>
+
+#define N 40
+
+int
+main ()
+{
+  int x[N];
+
+  #pragma omp target data map(x[:N]) use_device_addr(x)
+    #pragma omp target has_device_addr(x)
+      for (int i = 0; i < N; i++)
+	x[i] = i;
+  for (int i = 0; i < N; i++)
+    if (x[i] != i)
+      __builtin_abort ();
+
+  #pragma omp target data map(to: x[:N]) use_device_addr(x)
+    #pragma omp target has_device_addr(x)
+      for (int i = 0; i < N; i++)
+	x[i] = i+1;
+  for (int i = 0; i < N; i++)
+    if (x[i] != i)
+      __builtin_abort ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-has-device-addr-4.c b/libgomp/testsuite/libgomp.c/target-has-device-addr-4.c
new file mode 100644
index 0000000..9d89f38
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-has-device-addr-4.c
@@ -0,0 +1,34 @@
+/* Testing the 'has_device_addr' clause on the target construct with
+   enclosing 'target data' construct. */
+
+int
+foo (int size)
+{
+  int x[size];
+
+  #pragma omp target data map(x[:size]) use_device_addr(x)
+    #pragma omp target has_device_addr(x)
+      for (int i = 0; i < size; i++)
+	x[i] = i;
+  for (int i = 0; i < size; i++)
+    if (x[i] != i)
+      __builtin_abort ();
+
+  #pragma omp target data map(to: x[:size]) use_device_addr(x)
+    #pragma omp target has_device_addr(x)
+      for (int i = 0; i < size; i++)
+	x[i] = i+1;
+  for (int i = 0; i < size; i++)
+    if (x[i] != i)
+      __builtin_abort ();
+
+  return 0;
+}
+
+int
+main ()
+{
+  foo (40);
+
+  return 0;
+}
