public inbox for gcc-patches@gcc.gnu.org
 help / color / mirror / Atom feed
* [PATCH] C, C++, OpenMP: Add 'has_device_addr' clause to 'target' construct
@ 2021-10-18 16:17 Marcel Vollweiler
  2021-10-20 12:38 ` Jakub Jelinek
  0 siblings, 1 reply; 9+ messages in thread
From: Marcel Vollweiler @ 2021-10-18 16:17 UTC (permalink / raw)
  To: gcc-patches; +Cc: Jakub Jelinek

[-- Attachment #1: Type: text/plain, Size: 2166 bytes --]

Hi,

This patch adds the 'has_device_addr' clause to the OpenMP 'target'
construct which was introduced in OpenMP 5.1:

"The has_device_addr clause was added to the target construct to allow
access to variables or array sections that already have a device
address" (OpenMP 5.1 Specification, p. 669)

"The has_device_addr clause indicates that its list items already have
device addresses and therefore they may be directly accessed from a
target device. If the device address of a list item is not for the
device on which the target region executes, accessing the list item
inside the region results in unspecified behavior. The list items may
include array sections." (OpenMP 5.1 Specification, p. 200)

There are some restrictions for 'has_device_addr' (p. 202f):

1. "A list item may not be specified in both an is_device_ptr clause and
a has_device_addr clause on the directive."

2. "A list item that appears in an is_device_ptr or a has_device_addr
clause must not be specified in any data-sharing attribute clause on the
same target construct."

3. "A list item that appears in a has_device_addr clause must have a
valid device address for the device data environment."

4. As discussed on the omp-lang mailing list
(https://mailman.openmp.org/mailman/private/omp-lang/2021/017982.html),
has_device_addr is a data-sharing attribute clause (that is not yet
stated explicitly but will be corrected in OpenMP 5.2) and should not be
used together with the map clause on the same construct.

Similar restrictions hold also for the 'is_device_ptr' clause, so I
updated the code and added tests for that clause, too.

I tested the patch without regressions on powerpc64le-linux-gnu with
nvptx offloading and x86_64-linux-gnu with amdgcn offloading.

This patch only considers C/C++. The changes for Fortran will be
submitted separately later.

Marcel
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955

[-- Attachment #2: has-device-addr-patch.diff --]
[-- Type: text/plain, Size: 54881 bytes --]

C, C++, OpenMP: Add 'has_device_addr' clause to 'target' construct.

This patch adds the 'has_device_addr' clause to the OpenMP 'target' construct
which was introduced in OpenMP 5.1.

gcc/c-family/ChangeLog:

	* c-omp.c (c_omp_split_clauses): Add OMP_CLAUSE_HAS_DEVICE_ADDR case.
	* c-pragma.h (enum pragma_kind): Add 5.1 in comment.
	(enum pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR.

gcc/c/ChangeLog:

	* c-parser.c (c_parser_omp_clause_name): Parse 'has_device_addr' clause.
	(c_parser_omp_clause_has_device_addr): Added.
	(c_parser_omp_all_clauses): Add PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR case.
	(c_parser_omp_target_exit_data): Add HAS_DEVICE_ADDR to OMP_CLAUSE_MASK.
	* c-typeck.c (c_finish_omp_clauses): Add check that has_device_addr and 
	is_device_ptr do not appear together with map.

gcc/cp/ChangeLog:

	* parser.c (cp_parser_omp_clause_name): Parse 'has_device_addr' clause.
	(cp_parser_omp_all_clauses): Add PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR case.
	(cp_parser_omp_target_update): Add HAS_DEVICE_ADDR to OMP_CLAUSE_MASK.
	* pt.c (tsubst_omp_clauses): Add cases for OMP_CLAUSE_HAS_DEVICE_ADDR.
	* semantics.c (finish_omp_clauses): Add check that has_device_addr and
	is_device_ptr do not appear together with map.

gcc/ChangeLog:

	* gimplify.c (gimplify_scan_omp_clauses): Add OMP_CLAUSE_HAS_DEVICE_ADDR case.
	(gimplify_adjust_omp_clauses): Likewise.
	* omp-low.c (scan_sharing_clauses): Add lowering for has_device_addr clause.
	(lower_omp_target): Likewise.
	* tree-core.h (enum omp_clause_code): Update enum.
	* tree-nested.c (convert_nonlocal_omp_clauses): Add has_device_addr support.
	(convert_local_omp_clauses): Likewise.
	* tree-pretty-print.c (dump_omp_clause): Likewise.
	* tree.c: Update omp_clause_num_ops array.

libgomp/ChangeLog:

	* libgomp.texi: Updated entry for 'has-device-addr'.
	* testsuite/libgomp.c++/target-has-device-addr-2.C: New test.
	* testsuite/libgomp.c-c++-common/target-has-device-addr-1.c: New test.
	* testsuite/libgomp.c-c++-common/target-has-device-addr-3.c: New test.
	* testsuite/libgomp.c/target-has-device-addr-4.c: New test.

gcc/testsuite/ChangeLog:

	* c-c++-common/gomp/clauses-1.c: Add has_device_addr to test cases.
	* g++.dg/gomp/attrs-1.C: Likewise.
	* g++.dg/gomp/attrs-2.C: Likewise.
	* c-c++-common/gomp/target-has-device-addr-1.c: New test.
	* c-c++-common/gomp/target-is-device-ptr.c: New test.

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

^ permalink raw reply	[flat|nested] 9+ messages in thread

end of thread, other threads:[~2022-02-02 14:24 UTC | newest]

Thread overview: 9+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2021-10-18 16:17 [PATCH] C, C++, OpenMP: Add 'has_device_addr' clause to 'target' construct Marcel Vollweiler
2021-10-20 12:38 ` Jakub Jelinek
2021-11-15  9:03   ` Marcel Vollweiler
2021-11-24 17:08     ` [PATCH] C, C++, Fortran, " Marcel Vollweiler
2022-01-11 11:53       ` Jakub Jelinek
2022-01-11 13:27         ` Tobias Burnus
2022-01-11 14:07           ` Jakub Jelinek
2022-02-02  8:19         ` Marcel Vollweiler
2022-02-02 14:24           ` Jakub Jelinek

This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).