public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc r12-7161] C, C++, Fortran, OpenMP: Add 'has_device_addr' clause to 'target' construct.
@ 2022-02-10 7:50 Marcel Vollweiler
0 siblings, 0 replies; only message in thread
From: Marcel Vollweiler @ 2022-02-10 7:50 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:bbb7f8604e1dfc08f44354cfd93d2287f2fdd489
commit r12-7161-gbbb7f8604e1dfc08f44354cfd93d2287f2fdd489
Author: Marcel Vollweiler <marcel@codesourcery.com>
Date: Wed Feb 9 23:47:12 2022 -0800
C, C++, Fortran, 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 (OpenMP API 5.1 specification pp. 197ff):
has_device_addr(list)
"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." (p. 200)
"A list item may not be specified in both an is_device_ptr clause and a
has_device_addr clause on the directive." (p. 202)
"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." (p. 203)
gcc/c-family/ChangeLog:
* c-omp.cc (c_omp_split_clauses): Added OMP_CLAUSE_HAS_DEVICE_ADDR case.
* c-pragma.h (enum pragma_kind): Added 5.1 in comment.
(enum pragma_omp_clause): Added PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR.
gcc/c/ChangeLog:
* c-parser.cc (c_parser_omp_clause_name): Parse 'has_device_addr'
clause.
(c_parser_omp_variable_list): Handle array sections.
(c_parser_omp_clause_has_device_addr): Added.
(c_parser_omp_all_clauses): Added PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR
case.
(c_parser_omp_target_exit_data): Added HAS_DEVICE_ADDR to
OMP_CLAUSE_MASK.
* c-typeck.cc (handle_omp_array_sections): Handle clause restrictions.
(c_finish_omp_clauses): Handle array sections.
gcc/cp/ChangeLog:
* parser.cc (cp_parser_omp_clause_name): Parse 'has_device_addr' clause.
(cp_parser_omp_var_list_no_open): Handle array sections.
(cp_parser_omp_all_clauses): Added PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR
case.
(cp_parser_omp_target_update): Added HAS_DEVICE_ADDR to OMP_CLAUSE_MASK.
* semantics.cc (handle_omp_array_sections): Handle clause restrictions.
(finish_omp_clauses): Handle array sections.
gcc/fortran/ChangeLog:
* dump-parse-tree.cc (show_omp_clauses): Added OMP_LIST_HAS_DEVICE_ADDR
case.
* gfortran.h: Added OMP_LIST_HAS_DEVICE_ADDR.
* openmp.cc (enum omp_mask2): Added OMP_CLAUSE_HAS_DEVICE_ADDR.
(gfc_match_omp_clauses): Parse HAS_DEVICE_ADDR clause.
(resolve_omp_clauses): Same.
* trans-openmp.cc (gfc_trans_omp_variable_list): Added
OMP_LIST_HAS_DEVICE_ADDR case.
(gfc_trans_omp_clauses): Firstprivatize of array descriptors.
gcc/ChangeLog:
* gimplify.cc (gimplify_scan_omp_clauses): Added cases for
OMP_CLAUSE_HAS_DEVICE_ADDR
and handle array sections.
(gimplify_adjust_omp_clauses): Added OMP_CLAUSE_HAS_DEVICE_ADDR case.
* omp-low.cc (scan_sharing_clauses): Handle OMP_CLAUSE_HAS_DEVICE_ADDR.
(lower_omp_target): Same.
* tree-core.h (enum omp_clause_code): Same.
* tree-nested.cc (convert_nonlocal_omp_clauses): Same.
(convert_local_omp_clauses): Same.
* tree-pretty-print.cc (dump_omp_clause): Same.
* tree.cc: Same.
libgomp/ChangeLog:
* libgomp.texi: Updated entry for HAS_DEVICE_ADDR.
* target.c (copy_firstprivate_data): Copy only if host address is not
NULL.
* testsuite/libgomp.c++/target-has-device-addr-2.C: New test.
* testsuite/libgomp.c++/target-has-device-addr-4.C: New test.
* testsuite/libgomp.c++/target-has-device-addr-5.C: New test.
* testsuite/libgomp.c++/target-has-device-addr-6.C: New test.
* testsuite/libgomp.c-c++-common/target-has-device-addr-1.c: New test.
* testsuite/libgomp.c/target-has-device-addr-3.c: New test.
* testsuite/libgomp.fortran/target-has-device-addr-1.f90: New test.
* testsuite/libgomp.fortran/target-has-device-addr-2.f90: New test.
* testsuite/libgomp.fortran/target-has-device-addr-3.f90: New test.
* testsuite/libgomp.fortran/target-has-device-addr-4.f90: New test.
gcc/testsuite/ChangeLog:
* c-c++-common/gomp/clauses-1.c: Added has_device_addr to test cases.
* g++.dg/gomp/attrs-1.C: Added has_device_addr to test cases.
* g++.dg/gomp/attrs-2.C: Added has_device_addr to test cases.
* c-c++-common/gomp/target-has-device-addr-1.c: New test.
* c-c++-common/gomp/target-has-device-addr-2.c: New test.
* c-c++-common/gomp/target-is-device-ptr-1.c: New test.
* c-c++-common/gomp/target-is-device-ptr-2.c: New test.
* gfortran.dg/gomp/is_device_ptr-3.f90: New test.
* gfortran.dg/gomp/target-has-device-addr-1.f90: New test.
* gfortran.dg/gomp/target-has-device-addr-2.f90: New test.
Diff:
---
gcc/c-family/c-omp.cc | 1 +
gcc/c-family/c-pragma.h | 5 +-
gcc/c/c-parser.cc | 22 +++++-
gcc/c/c-typeck.cc | 32 +++++++-
gcc/cp/parser.cc | 13 +++-
gcc/cp/semantics.cc | 33 +++++++-
gcc/fortran/dump-parse-tree.cc | 1 +
gcc/fortran/gfortran.h | 3 +-
gcc/fortran/openmp.cc | 18 +++--
gcc/fortran/trans-openmp.cc | 15 +++-
gcc/gimplify.cc | 19 ++++-
gcc/omp-low.cc | 79 ++++++++++++++++---
gcc/testsuite/c-c++-common/gomp/clauses-1.c | 31 ++++----
.../c-c++-common/gomp/target-has-device-addr-1.c | 65 ++++++++++++++++
.../c-c++-common/gomp/target-has-device-addr-2.c | 17 ++++
.../c-c++-common/gomp/target-is-device-ptr-1.c | 22 ++++++
.../c-c++-common/gomp/target-is-device-ptr-2.c | 17 ++++
gcc/testsuite/g++.dg/gomp/attrs-1.C | 33 ++++----
gcc/testsuite/g++.dg/gomp/attrs-2.C | 35 +++++----
gcc/testsuite/gfortran.dg/gomp/is_device_ptr-3.f90 | 27 +++++++
.../gfortran.dg/gomp/target-has-device-addr-1.f90 | 36 +++++++++
.../gfortran.dg/gomp/target-has-device-addr-2.f90 | 27 +++++++
gcc/tree-core.h | 3 +
gcc/tree-nested.cc | 2 +
gcc/tree-pretty-print.cc | 3 +
gcc/tree.cc | 2 +
libgomp/libgomp.texi | 2 +-
libgomp/target.c | 2 +-
.../libgomp.c++/target-has-device-addr-2.C | 23 ++++++
.../libgomp.c++/target-has-device-addr-4.C | 33 ++++++++
.../libgomp.c++/target-has-device-addr-5.C | 33 ++++++++
.../libgomp.c++/target-has-device-addr-6.C | 32 ++++++++
.../target-has-device-addr-1.c | 73 ++++++++++++++++++
.../testsuite/libgomp.c/target-has-device-addr-3.c | 33 ++++++++
.../libgomp.fortran/target-has-device-addr-1.f90 | 50 ++++++++++++
.../libgomp.fortran/target-has-device-addr-2.f90 | 40 ++++++++++
.../libgomp.fortran/target-has-device-addr-3.f90 | 90 ++++++++++++++++++++++
.../libgomp.fortran/target-has-device-addr-4.f90 | 71 +++++++++++++++++
38 files changed, 961 insertions(+), 82 deletions(-)
diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index 0251aec5348..f5314d66f27 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -1862,6 +1862,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 6e6c806ccf7..54864c2ec41 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.cc b/gcc/c/c-parser.cc
index 7e81c33de11..3b1d2d45add 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -12771,7 +12771,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;
@@ -13164,6 +13166,7 @@ c_parser_omp_variable_list (c_parser *parser,
case OMP_CLAUSE_REDUCTION:
case OMP_CLAUSE_IN_REDUCTION:
case OMP_CLAUSE_TASK_REDUCTION:
+ case OMP_CLAUSE_HAS_DEVICE_ADDR:
array_section_p = false;
dims.truncate (0);
while (c_parser_next_token_is (parser, CPP_OPEN_SQUARE))
@@ -14324,6 +14327,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 ) */
@@ -17052,6 +17065,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";
@@ -21034,7 +21051,8 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEFAULTMAP) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IN_REDUCTION) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_THREAD_LIMIT) \
- | (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.cc b/gcc/c/c-typeck.cc
index 39094cc6f50..3075c883548 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -13804,6 +13804,8 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
}
first = c_fully_fold (first, false, NULL);
OMP_CLAUSE_DECL (c) = first;
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
+ return false;
if (size)
size = c_fully_fold (size, false, NULL);
OMP_CLAUSE_SIZE (c) = size;
@@ -14109,7 +14111,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, is_on_device_head;
tree c, t, type, *pc;
tree simdlen = NULL_TREE, safelen = NULL_TREE;
bool branch_seen = false;
@@ -14145,6 +14147,7 @@ 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 (&is_on_device_head, &bitmap_default_obstack);
if (ort & C_ORT_ACC)
for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
@@ -14573,7 +14576,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)
@@ -15187,7 +15192,8 @@ 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 (&is_on_device_head, DECL_UID (t)))
{
if (ort == C_ORT_ACC)
error_at (OMP_CLAUSE_LOCATION (c),
@@ -15272,6 +15278,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_on_device_head, DECL_UID (t));
if (TREE_CODE (TREE_TYPE (t)) != POINTER_TYPE)
{
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR
@@ -15292,6 +15300,24 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
goto check_dup_generic;
+ case OMP_CLAUSE_HAS_DEVICE_ADDR:
+ t = OMP_CLAUSE_DECL (c);
+ if (TREE_CODE (t) == TREE_LIST)
+ {
+ if (handle_omp_array_sections (c, ort))
+ remove = true;
+ else
+ {
+ t = OMP_CLAUSE_DECL (c);
+ while (TREE_CODE (t) == ARRAY_REF)
+ t = TREE_OPERAND (t, 0);
+ }
+ }
+ bitmap_set_bit (&is_on_device_head, DECL_UID (t));
+ if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
+ c_mark_addressable (t);
+ goto check_dup_generic_t;
+
case OMP_CLAUSE_USE_DEVICE_ADDR:
t = OMP_CLAUSE_DECL (c);
if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 94a5c64be4c..03d99aba13e 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -36341,7 +36341,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;
@@ -36644,6 +36646,7 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
case OMP_CLAUSE_REDUCTION:
case OMP_CLAUSE_IN_REDUCTION:
case OMP_CLAUSE_TASK_REDUCTION:
+ case OMP_CLAUSE_HAS_DEVICE_ADDR:
array_section_p = false;
dims.truncate (0);
while (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_SQUARE))
@@ -40085,6 +40088,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);
@@ -44265,7 +44273,8 @@ cp_parser_omp_target_update (cp_parser *parser, cp_token *pragma_tok,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_ALLOCATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IN_REDUCTION) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_THREAD_LIMIT) \
- | (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/semantics.cc b/gcc/cp/semantics.cc
index 466d6b56871..0cb17a6a8ab 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -5648,6 +5648,8 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
return false;
}
OMP_CLAUSE_DECL (c) = first;
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
+ return false;
OMP_CLAUSE_SIZE (c) = size;
if (TREE_CODE (t) == FIELD_DECL)
t = finish_non_static_data_member (t, NULL_TREE, NULL_TREE);
@@ -6677,7 +6679,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, is_on_device_head;
tree c, t, *pc;
tree safelen = NULL_TREE;
bool branch_seen = false;
@@ -6710,6 +6712,7 @@ 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 (&is_on_device_head, &bitmap_default_obstack);
if (ort & C_ORT_ACC)
for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
@@ -7008,7 +7011,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)
@@ -8232,7 +8237,8 @@ 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 (&is_on_device_head, DECL_UID (t)))
{
if (ort == C_ORT_ACC)
error_at (OMP_CLAUSE_LOCATION (c),
@@ -8491,6 +8497,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_on_device_head, DECL_UID (t));
if (!type_dependent_expression_p (t))
{
tree type = TREE_TYPE (t);
@@ -8520,6 +8528,25 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
goto check_dup_generic;
+ case OMP_CLAUSE_HAS_DEVICE_ADDR:
+ t = OMP_CLAUSE_DECL (c);
+ if (TREE_CODE (t) == TREE_LIST)
+ {
+ if (handle_omp_array_sections (c, ort))
+ remove = true;
+ else
+ {
+ t = OMP_CLAUSE_DECL (c);
+ while (TREE_CODE (t) == INDIRECT_REF
+ || TREE_CODE (t) == ARRAY_REF)
+ t = TREE_OPERAND (t, 0);
+ }
+ }
+ bitmap_set_bit (&is_on_device_head, DECL_UID (t));
+ if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
+ cxx_mark_addressable (t);
+ goto check_dup_generic_t;
+
case OMP_CLAUSE_USE_DEVICE_ADDR:
field_ok = true;
t = OMP_CLAUSE_DECL (c);
diff --git a/gcc/fortran/dump-parse-tree.cc b/gcc/fortran/dump-parse-tree.cc
index 3112caec053..2a2f9901b08 100644
--- a/gcc/fortran/dump-parse-tree.cc
+++ b/gcc/fortran/dump-parse-tree.cc
@@ -1683,6 +1683,7 @@ show_omp_clauses (gfc_omp_clauses *omp_clauses)
case OMP_LIST_CACHE: type = "CACHE"; break;
case OMP_LIST_IS_DEVICE_PTR: type = "IS_DEVICE_PTR"; break;
case OMP_LIST_USE_DEVICE_PTR: type = "USE_DEVICE_PTR"; break;
+ case OMP_LIST_HAS_DEVICE_ADDR: type = "HAS_DEVICE_ADDR"; break;
case OMP_LIST_USE_DEVICE_ADDR: type = "USE_DEVICE_ADDR"; break;
case OMP_LIST_NONTEMPORAL: type = "NONTEMPORAL"; break;
case OMP_LIST_ALLOCATE: type = "ALLOCATE"; break;
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index 993879feda4..cb136f875f4 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1393,7 +1393,8 @@ enum
OMP_LIST_USE_DEVICE_ADDR,
OMP_LIST_NONTEMPORAL,
OMP_LIST_ALLOCATE,
- OMP_LIST_NUM
+ OMP_LIST_HAS_DEVICE_ADDR,
+ OMP_LIST_NUM /* Must be the last. */
};
/* Because a symbol can belong to multiple namelists, they must be
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index 38c67e1f640..33b372f3430 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -926,7 +926,7 @@ enum omp_mask1
OMP_MASK1_LAST
};
-/* OpenACC 2.0+ specific clauses. */
+/* More OpenMP clauses and OpenACC 2.0+ specific clauses. */
enum omp_mask2
{
OMP_CLAUSE_ASYNC,
@@ -955,6 +955,7 @@ enum omp_mask2
OMP_CLAUSE_FINALIZE,
OMP_CLAUSE_ATTACH,
OMP_CLAUSE_NOHOST,
+ OMP_CLAUSE_HAS_DEVICE_ADDR, /* OpenMP 5.1 */
/* This must come last. */
OMP_MASK2_LAST
};
@@ -2151,6 +2152,11 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
}
break;
case 'h':
+ if ((mask & OMP_CLAUSE_HAS_DEVICE_ADDR)
+ && gfc_match_omp_variable_list
+ ("has_device_addr (", &c->lists[OMP_LIST_HAS_DEVICE_ADDR],
+ false, NULL, NULL, true) == MATCH_YES)
+ continue;
if ((mask & OMP_CLAUSE_HINT)
&& (m = gfc_match_dupl_check (!c->hint, "hint", true, &c->hint))
!= MATCH_NO)
@@ -2923,8 +2929,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
continue;
if ((mask & OMP_CLAUSE_USE_DEVICE_ADDR)
&& gfc_match_omp_variable_list
- ("use_device_addr (",
- &c->lists[OMP_LIST_USE_DEVICE_ADDR], false) == MATCH_YES)
+ ("use_device_addr (", &c->lists[OMP_LIST_USE_DEVICE_ADDR],
+ false, NULL, NULL, true) == MATCH_YES)
continue;
break;
case 'v':
@@ -3651,7 +3657,8 @@ cleanup:
| OMP_CLAUSE_DEPEND | OMP_CLAUSE_NOWAIT | OMP_CLAUSE_PRIVATE \
| OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULTMAP \
| OMP_CLAUSE_IS_DEVICE_PTR | OMP_CLAUSE_IN_REDUCTION \
- | OMP_CLAUSE_THREAD_LIMIT | OMP_CLAUSE_ALLOCATE)
+ | OMP_CLAUSE_THREAD_LIMIT | OMP_CLAUSE_ALLOCATE \
+ | OMP_CLAUSE_HAS_DEVICE_ADDR)
#define OMP_TARGET_DATA_CLAUSES \
(omp_mask (OMP_CLAUSE_DEVICE) | OMP_CLAUSE_MAP | OMP_CLAUSE_IF \
| OMP_CLAUSE_USE_DEVICE_PTR | OMP_CLAUSE_USE_DEVICE_ADDR)
@@ -6283,7 +6290,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
"IN_REDUCTION", "TASK_REDUCTION",
"DEVICE_RESIDENT", "LINK", "USE_DEVICE",
"CACHE", "IS_DEVICE_PTR", "USE_DEVICE_PTR", "USE_DEVICE_ADDR",
- "NONTEMPORAL", "ALLOCATE" };
+ "NONTEMPORAL", "ALLOCATE", "HAS_DEVICE_ADDR" };
STATIC_ASSERT (ARRAY_SIZE (clause_names) == OMP_LIST_NUM);
if (omp_clauses == NULL)
@@ -7132,6 +7139,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
n->sym->name, name, &n->where);
}
break;
+ case OMP_LIST_HAS_DEVICE_ADDR:
case OMP_LIST_USE_DEVICE_PTR:
case OMP_LIST_USE_DEVICE_ADDR:
/* FIXME: Handle OMP_LIST_USE_DEVICE_PTR. */
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index d5a6b2d6ee3..0eba0b3c3e1 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -1910,7 +1910,17 @@ gfc_trans_omp_variable_list (enum omp_clause_code code,
tree t = gfc_trans_omp_variable (namelist->sym, declare_simd);
if (t != error_mark_node)
{
- tree node = build_omp_clause (input_location, code);
+ tree node;
+ /* For HAS_DEVICE_ADDR of an array descriptor, firstprivatize the
+ descriptor such that the bounds are available; its data component
+ is unmodified; it is handled as device address inside target. */
+ if (code == OMP_CLAUSE_HAS_DEVICE_ADDR
+ && (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (t))
+ || (POINTER_TYPE_P (TREE_TYPE (t))
+ && GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (TREE_TYPE (t))))))
+ node = build_omp_clause (input_location, OMP_CLAUSE_FIRSTPRIVATE);
+ else
+ node = build_omp_clause (input_location, code);
OMP_CLAUSE_DECL (node) = t;
list = gfc_trans_add_clause (node, list);
@@ -2604,6 +2614,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
case OMP_LIST_IS_DEVICE_PTR:
clause_code = OMP_CLAUSE_IS_DEVICE_PTR;
goto add_clause;
+ case OMP_LIST_HAS_DEVICE_ADDR:
+ clause_code = OMP_CLAUSE_HAS_DEVICE_ADDR;
+ goto add_clause;
case OMP_LIST_NONTEMPORAL:
clause_code = OMP_CLAUSE_NONTEMPORAL;
goto add_clause;
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 875b115d02d..8d676fb96c8 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -10278,6 +10278,14 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
flags = GOVD_EXPLICIT;
goto do_add;
+ case OMP_CLAUSE_HAS_DEVICE_ADDR:
+ decl = OMP_CLAUSE_DECL (c);
+ while (TREE_CODE (decl) == INDIRECT_REF
+ || TREE_CODE (decl) == ARRAY_REF)
+ decl = TREE_OPERAND (decl, 0);
+ flags = GOVD_EXPLICIT;
+ goto do_add_decl;
+
case OMP_CLAUSE_IS_DEVICE_PTR:
flags = GOVD_FIRSTPRIVATE | GOVD_EXPLICIT;
goto do_add;
@@ -11428,6 +11436,16 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
}
break;
+ case OMP_CLAUSE_HAS_DEVICE_ADDR:
+ decl = OMP_CLAUSE_DECL (c);
+ while (TREE_CODE (decl) == INDIRECT_REF
+ || TREE_CODE (decl) == ARRAY_REF)
+ decl = TREE_OPERAND (decl, 0);
+ n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
+ remove = n == NULL || !(n->value & GOVD_SEEN);
+ break;
+
+ case OMP_CLAUSE_IS_DEVICE_PTR:
case OMP_CLAUSE_NONTEMPORAL:
decl = OMP_CLAUSE_DECL (c);
n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
@@ -11729,7 +11747,6 @@ 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_IS_DEVICE_PTR:
case OMP_CLAUSE_ASYNC:
case OMP_CLAUSE_WAIT:
case OMP_CLAUSE_INDEPENDENT:
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index c33b3daa439..065208464b2 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -1375,7 +1375,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)
@@ -1383,8 +1384,14 @@ 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 (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
+ {
+ if (TREE_CODE (decl) == INDIRECT_REF)
+ decl = TREE_OPERAND (decl, 0);
+ install_var_field (decl, true, 3, ctx);
+ }
else if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
- install_var_field (decl, true, 3, ctx);
+ install_var_field (decl, true, 3, ctx);
else
install_var_field (decl, false, 3, ctx);
}
@@ -1452,6 +1459,13 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
install_var_local (decl, ctx);
break;
+ case OMP_CLAUSE_HAS_DEVICE_ADDR:
+ decl = OMP_CLAUSE_DECL (c);
+ while (TREE_CODE (decl) == INDIRECT_REF
+ || TREE_CODE (decl) == ARRAY_REF)
+ decl = TREE_OPERAND (decl, 0);
+ goto do_private;
+
case OMP_CLAUSE_IS_DEVICE_PTR:
decl = OMP_CLAUSE_DECL (c);
goto do_private;
@@ -1729,12 +1743,21 @@ 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 (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
+ {
+ while (TREE_CODE (decl) == INDIRECT_REF
+ || TREE_CODE (decl) == ARRAY_REF)
+ decl = TREE_OPERAND (decl, 0);
+ }
+
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);
@@ -12819,8 +12842,15 @@ 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);
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
+ {
+ while (TREE_CODE (var) == INDIRECT_REF
+ || TREE_CODE (var) == ARRAY_REF)
+ var = TREE_OPERAND (var, 0);
+ }
map_cnt++;
if (is_variable_sized (var))
{
@@ -12835,7 +12865,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))
@@ -13301,17 +13332,26 @@ 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);
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
+ {
+ while (TREE_CODE (ovar) == INDIRECT_REF
+ || TREE_CODE (ovar) == ARRAY_REF)
+ ovar = TREE_OPERAND (ovar, 0);
+ }
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);
@@ -13333,7 +13373,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)
@@ -13348,6 +13389,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))))
@@ -13545,6 +13587,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;
@@ -13555,12 +13598,21 @@ 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);
else
- x = build_receiver_ref (var, false, ctx);
+ {
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
+ {
+ while (TREE_CODE (var) == INDIRECT_REF
+ || TREE_CODE (var) == ARRAY_REF)
+ var = TREE_OPERAND (var, 0);
+ }
+ x = build_receiver_ref (var, false, ctx);
+ }
if (is_array_data)
{
@@ -13607,7 +13659,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)
@@ -13630,7 +13683,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
type = TREE_TYPE (type);
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_USE_DEVICE_ADDR
+ && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_HAS_DEVICE_ADDR)
|| (omp_privatize_by_reference (var)
&& omp_is_allocatable_or_ptr (var))))
{
@@ -13653,7 +13707,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
gimple_build_assign (new_var, x));
}
tree present;
- present = (do_optional_check
+ present = ((do_optional_check
+ && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_HAS_DEVICE_ADDR)
? omp_check_optional_argument (OMP_CLAUSE_DECL (c), true)
: NULL_TREE);
if (present)
diff --git a/gcc/testsuite/c-c++-common/gomp/clauses-1.c b/gcc/testsuite/c-c++-common/gomp/clauses-1.c
index 3ff49e0a298..71ca41c5804 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 - 1: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-1: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-1: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 00000000000..ebf55eea70b
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-has-device-addr-1.c
@@ -0,0 +1,65 @@
+/* { 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" } */
+ ;
+
+ int z[3] = { 2, 5, 7 };
+ #pragma omp target data map(z[:3]) use_device_addr(z)
+ #pragma omp target has_device_addr(z[1:])
+ ;
+
+ #pragma omp target data map(z[:3]) use_device_addr(z)
+ #pragma omp target has_device_addr(z[1])
+ ;
+
+ #pragma omp target data map(z[:3]) use_device_addr(z)
+ #pragma omp target has_device_addr(z[1:2])
+ ;
+
+ #pragma omp target data map(z[:3]) use_device_addr(z)
+ #pragma omp target has_device_addr(z[:2])
+ ;
+
+ int w[3][4];
+ #pragma omp target data map(w) use_device_addr(w)
+ #pragma omp target has_device_addr(w[1][2])
+ ;
+
+ #pragma omp target data map(w) use_device_addr(w)
+ #pragma omp target has_device_addr(w[:1][2:])
+ ;
+
+ int u[0];
+ #pragma omp target data map(u) use_device_addr(u)
+ #pragma omp target has_device_addr(u)
+ ;
+
+ struct S { int m; } s;
+ s.m = 42;
+ #pragma omp target data map (s) use_device_addr (s)
+ #pragma omp target has_device_addr (s)
+ ++s.m;
+
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/target-has-device-addr-2.c b/gcc/testsuite/c-c++-common/gomp/target-has-device-addr-2.c
new file mode 100644
index 00000000000..7378416964c
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-has-device-addr-2.c
@@ -0,0 +1,17 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -fdump-tree-gimple" } */
+
+void
+foo ()
+{
+ int x, y;
+
+ #pragma omp target data map(x, y) use_device_addr(x, y)
+ #pragma omp target has_device_addr(x, y)
+ {
+ x = 42;
+ }
+}
+
+/* { dg-final { scan-tree-dump "has_device_addr\\(x\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump-not "has_device_addr\\(y\\)" "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/target-is-device-ptr-1.c b/gcc/testsuite/c-c++-common/gomp/target-is-device-ptr-1.c
new file mode 100644
index 00000000000..ecf30ca65f7
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-is-device-ptr-1.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/c-c++-common/gomp/target-is-device-ptr-2.c b/gcc/testsuite/c-c++-common/gomp/target-is-device-ptr-2.c
new file mode 100644
index 00000000000..df743dda166
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-is-device-ptr-2.c
@@ -0,0 +1,17 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -fdump-tree-gimple" } */
+
+void
+foo ()
+{
+ int *x, *y;
+
+ #pragma omp target data map(x, y) use_device_ptr(x, y)
+ #pragma omp target is_device_ptr(x, y)
+ {
+ *x = 42;
+ }
+}
+
+/* { dg-final { scan-tree-dump "is_device_ptr\\(x\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump-not "is_device_ptr\\(y\\)" "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/attrs-1.C b/gcc/testsuite/g++.dg/gomp/attrs-1.C
index 319ad3241de..f64b078db40 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-1: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-1: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 - 1 : 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 955b2dd04c7..cc91fa28307 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-1: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-1: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/gfortran.dg/gomp/is_device_ptr-3.f90 b/gcc/testsuite/gfortran.dg/gomp/is_device_ptr-3.f90
new file mode 100644
index 00000000000..c3de7726e88
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/is_device_ptr-3.f90
@@ -0,0 +1,27 @@
+! Test to ensure that IS_DEVICE_PTR is removed for non-used variables.
+
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+program main
+ use iso_c_binding
+ implicit none
+
+ integer :: x, y
+ call foo (x, y)
+
+contains
+ subroutine foo (a, b)
+ integer, target :: a, b
+
+ !$omp target data map(a, b) use_device_ptr(a, b)
+ !$omp target is_device_ptr(a, b)
+ a = 42
+ !$omp end target
+ !$omp end target data
+ end subroutine foo
+
+end program main
+
+! { dg-final { scan-tree-dump "is_device_ptr\\(a\\)" "gimple" } }
+! { dg-final { scan-tree-dump-not "is_device_ptr\\(b\\)" "gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-has-device-addr-1.f90 b/gcc/testsuite/gfortran.dg/gomp/target-has-device-addr-1.f90
new file mode 100644
index 00000000000..db3fa46b4a1
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target-has-device-addr-1.f90
@@ -0,0 +1,36 @@
+! { dg-do compile }
+
+implicit none
+
+integer, target :: x
+integer, pointer :: ptr
+integer :: a(5)
+
+!$omp target has_device_addr(x)
+!$omp end target
+!$omp target has_device_addr(ptr)
+!$omp end target
+!$omp target has_device_addr(a)
+!$omp end target
+!$omp target has_device_addr(a(2:3))
+!$omp end target
+!$omp target has_device_addr(a(:3))
+!$omp end target
+!$omp target has_device_addr(a(2:))
+!$omp end target
+!$omp target has_device_addr(a(2))
+!$omp end target
+
+!$omp target has_device_addr(x) has_device_addr(x) ! { dg-error "'x' present on multiple clauses" }
+!$omp end target
+
+!$omp target private(x) has_device_addr(x) ! { dg-error "'x' present on multiple clauses" }
+!$omp end target
+!$omp target has_device_addr(x) private(x) ! { dg-error "'x' present on multiple clauses" }
+!$omp end target
+!$omp target firstprivate(x) has_device_addr(x) ! { dg-error "'x' present on multiple clauses" }
+!$omp end target
+!$omp target has_device_addr(x) firstprivate(x) ! { dg-error "'x' present on multiple clauses" }
+!$omp end target
+
+end
\ No newline at end of file
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-has-device-addr-2.f90 b/gcc/testsuite/gfortran.dg/gomp/target-has-device-addr-2.f90
new file mode 100644
index 00000000000..7fc92b3cb6e
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target-has-device-addr-2.f90
@@ -0,0 +1,27 @@
+! Test to ensure that HAS_DEVICE_ADDR is removed for non-used variables.
+
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+program main
+ use iso_c_binding
+ implicit none
+
+ integer :: x, y
+ call foo (x, y)
+
+contains
+ subroutine foo (a, b)
+ integer :: a, b
+
+ !$omp target data map(a) use_device_addr(a)
+ !$omp target has_device_addr(a)
+ a = 42
+ !$omp end target
+ !$omp end target data
+ end subroutine foo
+
+end program main
+
+! { dg-final { scan-tree-dump "has_device_addr\\(a\\)" "gimple" } }
+! { dg-final { scan-tree-dump-not "has_device_addr\\(b\\)" "gimple" } }
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index bf2efa61330..01a1ce499da 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -342,6 +342,9 @@ enum omp_clause_code {
OpenMP clause: map ({alloc:,to:,from:,tofrom:,}variable-list). */
OMP_CLAUSE_MAP,
+ /* OpenMP clause: has_device_addr (variable-list). */
+ OMP_CLAUSE_HAS_DEVICE_ADDR,
+
/* Internal structure to hold OpenACC cache directive's variable-list.
#pragma acc cache (variable-list). */
OMP_CLAUSE__CACHE_,
diff --git a/gcc/tree-nested.cc b/gcc/tree-nested.cc
index b7e9a3b472f..078ceab3ca3 100644
--- a/gcc/tree-nested.cc
+++ b/gcc/tree-nested.cc
@@ -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.cc b/gcc/tree-pretty-print.cc
index 666b7a70ea2..99af977979d 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -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.cc b/gcc/tree.cc
index dfcdf6822f1..aeb3b4cebcf 100644
--- a/gcc/tree.cc
+++ b/gcc/tree.cc
@@ -289,6 +289,7 @@ unsigned const char omp_clause_num_ops[] =
2, /* OMP_CLAUSE_FROM */
2, /* OMP_CLAUSE_TO */
2, /* OMP_CLAUSE_MAP */
+ 1, /* OMP_CLAUSE_HAS_DEVICE_ADDR */
2, /* OMP_CLAUSE__CACHE_ */
2, /* OMP_CLAUSE_GANG */
1, /* OMP_CLAUSE_ASYNC */
@@ -378,6 +379,7 @@ const char * const omp_clause_code_name[] =
"from",
"to",
"map",
+ "has_device_addr",
"_cache_",
"gang",
"async",
diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
index 424459f4442..161a423ac7c 100644
--- a/libgomp/libgomp.texi
+++ b/libgomp/libgomp.texi
@@ -294,7 +294,7 @@ 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 Y @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 Y @tab
@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/target.c b/libgomp/target.c
index 698ff14a05f..9017458885e 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -2510,7 +2510,7 @@ copy_firstprivate_data (char *tgt, size_t mapnum, void **hostaddrs,
tgt_size = 0;
size_t i;
for (i = 0; i < mapnum; i++)
- if ((kinds[i] & 0xff) == GOMP_MAP_FIRSTPRIVATE)
+ if ((kinds[i] & 0xff) == GOMP_MAP_FIRSTPRIVATE && hostaddrs[i] != NULL)
{
size_t align = (size_t) 1 << (kinds[i] >> 8);
tgt_size = (tgt_size + align - 1) & ~(align - 1);
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 00000000000..d9a309d7af4
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-has-device-addr-2.C
@@ -0,0 +1,23 @@
+/* Testing 'has_device_addr' clause on the target construct with reference. */
+
+#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++/target-has-device-addr-4.C b/libgomp/testsuite/libgomp.c++/target-has-device-addr-4.C
new file mode 100644
index 00000000000..6468c6c8433
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-has-device-addr-4.C
@@ -0,0 +1,33 @@
+#include <omp.h>
+
+int
+main ()
+{
+ int *dp = (int*)omp_target_alloc (30*sizeof(int), 0);
+
+ #pragma omp target is_device_ptr(dp)
+ for (int i = 0; i < 30; i++)
+ dp[i] = i;
+
+ int (&x)[30] = *static_cast<int(*)[30]>(static_cast<void*>(dp));
+
+ #pragma omp target has_device_addr(x)
+ for (int i = 0; i < 30; i++)
+ x[i] = 2 * i;
+
+ #pragma omp target has_device_addr(x)
+ for (int i = 0; i < 30; i++)
+ if (x[i] != 2 * i)
+ __builtin_abort ();
+
+ #pragma omp target has_device_addr(x[1:5])
+ for (int i = 1; i < 6; i++)
+ x[i] = 3 * i;
+
+ #pragma omp target has_device_addr(x[1:5])
+ for (int i = 1; i < 6; i++)
+ if (x[i] != 3 * i)
+ __builtin_abort ();
+
+ omp_target_free (dp, 0);
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-has-device-addr-5.C b/libgomp/testsuite/libgomp.c++/target-has-device-addr-5.C
new file mode 100644
index 00000000000..e847cdceb44
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-has-device-addr-5.C
@@ -0,0 +1,33 @@
+/* Testing 'has_device_addr' clause on the target construct with reference. */
+
+#include <omp.h>
+
+int
+main ()
+{
+ int *dpx = (int*)omp_target_alloc (sizeof(int), 0);
+ int **dpy = (int**)omp_target_alloc (sizeof(int*), 0);
+
+ #pragma omp target is_device_ptr(dpx, dpy)
+ {
+ *dpx = 42;
+ int z = 77;
+ *dpy = &z;
+ }
+
+ int& x = *dpx;
+ int*& y = *dpy;
+
+ #pragma omp target has_device_addr(x, y)
+ {
+ x = 24;
+ y = &x;
+ }
+
+ #pragma omp target has_device_addr(x, y)
+ if (x != 24 || y != &x)
+ __builtin_abort ();
+
+ omp_target_free(dpx, 0);
+ omp_target_free(dpy, 0);
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-has-device-addr-6.C b/libgomp/testsuite/libgomp.c++/target-has-device-addr-6.C
new file mode 100644
index 00000000000..141edb14dec
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-has-device-addr-6.C
@@ -0,0 +1,32 @@
+/* Testing 'has_device_addr' clause on the target construct with reference. */
+
+#include <omp.h>
+
+int
+main ()
+{
+ int *dpx = (int*)omp_target_alloc (sizeof(int), 0);
+ double *dpy = (double*)omp_target_alloc (sizeof(double), 0);
+
+ #pragma omp target is_device_ptr(dpx, dpy)
+ {
+ *dpx = 42;
+ *dpy = 43.5;
+ }
+
+ int &x = *dpx;
+ double &y = *dpy;
+
+ #pragma omp target has_device_addr(x, y)
+ {
+ x = 24;
+ y = 25.7;
+ }
+
+ #pragma omp target has_device_addr(y, x)
+ if (x != 24 || y != 25.7)
+ __builtin_abort ();
+
+ omp_target_free(dpx, 0);
+ omp_target_free(dpy, 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 00000000000..fcc5c9e8553
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-has-device-addr-1.c
@@ -0,0 +1,73 @@
+/* Testing the 'has_device_addr' clause on the target construct with
+ enclosing 'target data' construct. */
+
+#define N 40
+
+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 ();
+
+ int y[N];
+
+ for (int i = 0; i < N; i++)
+ y[i] = 42;
+ #pragma omp target data map(y) use_device_addr(y)
+ #pragma omp target has_device_addr(y)
+ for (int i = 0; i < N; i++)
+ y[i] = i;
+ for (int i = 0; i < N; i++)
+ if (y[i] != i)
+ __builtin_abort ();
+
+ #pragma omp target data map(y[:N]) use_device_addr(y)
+ #pragma omp target has_device_addr(y[:N])
+ for (int i = 0; i < N; i++)
+ y[i] = i + 2;
+ for (int i = 0; i < N; i++)
+ if (y[i] != i + 2)
+ __builtin_abort ();
+
+ #pragma omp target data map(y[:N]) use_device_addr(y)
+ #pragma omp target has_device_addr(y[24])
+ y[24] = 42;
+ if (y[24] != 42)
+ __builtin_abort ();
+
+ #pragma omp target data map(y[:N]) use_device_addr(y)
+ #pragma omp target has_device_addr(y[24:])
+ for (int i = 24; i < N; i++)
+ y[i] = i + 3;
+ for (int i = 24; i < N; i++)
+ if (y[i] != i + 3)
+ __builtin_abort ();
+
+ #pragma omp target data map(y[:N]) use_device_addr(y)
+ #pragma omp target has_device_addr(y[12:24])
+ for (int i = 12; i < 24; i++)
+ y[i] = i + 4;
+ for (int i = 12; i < 24; i++)
+ if (y[i] != i + 4)
+ __builtin_abort ();
+
+ int u[0];
+ #pragma omp target data map(u) use_device_addr(u)
+ #pragma omp target has_device_addr(u)
+ ;
+
+ struct S { int m; } s;
+ s.m = 42;
+ #pragma omp target data map (s) use_device_addr (s)
+ #pragma omp target has_device_addr (s)
+ ++s.m;
+ if (s.m != 43)
+ __builtin_abort ();
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c/target-has-device-addr-3.c b/libgomp/testsuite/libgomp.c/target-has-device-addr-3.c
new file mode 100644
index 00000000000..fd99a82f66a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-has-device-addr-3.c
@@ -0,0 +1,33 @@
+/* Testing 'has_device_addr' clause with variable sized array. */
+
+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(x) use_device_addr(x)
+ #pragma omp target has_device_addr(x[2:3])
+ for (int i = 0; i < size; i++)
+ x[i] = i;
+ for (int i = 0; i < size; i++)
+ if (x[i] != i)
+ __builtin_abort ();
+
+ return 0;
+}
+
+int
+main ()
+{
+ foo (40);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.fortran/target-has-device-addr-1.f90 b/libgomp/testsuite/libgomp.fortran/target-has-device-addr-1.f90
new file mode 100644
index 00000000000..2945864fa53
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-has-device-addr-1.f90
@@ -0,0 +1,50 @@
+program main
+ use omp_lib
+ use iso_c_binding
+ implicit none
+
+ integer, parameter :: N = 40
+ integer :: x, i
+ integer :: y (N)
+ integer :: u (0)
+
+ x = 24
+ !$omp target data map(x) use_device_addr(x)
+ !$omp target has_device_addr(x)
+ x = 42;
+ !$omp end target
+ !$omp end target data
+ if (x /= 42) stop 1
+
+ y = 42
+ !$omp target data map(y) use_device_addr(y)
+ !$omp target has_device_addr(y)
+ y = [(i, i=1, N)]
+ !$omp end target
+ !$omp end target data
+ if (any (y /= [(i, i = 1, N)])) stop 2
+
+ !$omp target data map(y(:N)) use_device_addr(y)
+ !$omp target has_device_addr(y(:N))
+ y = [(i+2, i=1, N)]
+ !$omp end target
+ !$omp end target data
+ if (any (y /= [(i+2, i = 1, N)])) stop 3
+
+ !$omp target data map(y) use_device_addr(y)
+ !$omp target has_device_addr(y(24:))
+ do i = 24, N
+ y(i) = i + 3
+ end do
+ !$omp end target
+ !$omp end target data
+ do i = 24, N
+ if (y(i) /= i + 3) stop 5
+ end do
+
+ !$omp target data map(u) use_device_addr(u)
+ !$omp target has_device_addr(u)
+ !$omp end target
+ !$omp end target data
+
+end program main
diff --git a/libgomp/testsuite/libgomp.fortran/target-has-device-addr-2.f90 b/libgomp/testsuite/libgomp.fortran/target-has-device-addr-2.f90
new file mode 100644
index 00000000000..a8d78a75af3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-has-device-addr-2.f90
@@ -0,0 +1,40 @@
+program main
+ use omp_lib
+ use iso_c_binding
+ implicit none
+
+ integer, parameter :: N = 5
+ integer :: i, x(N), y(N), z(N:2*N-1)
+ target :: z
+
+ x = 42
+ y = 43
+ z = 44
+
+ call foo (x, y, z)
+ if (any (x /= [(i, i = 1, N)])) stop 1
+ if (any (y /= [(2*i, i = 1, N)])) stop 2
+ if (any (z /= [(3*i, i = 1, N)])) stop 3
+
+ contains
+ subroutine foo(a, b, c)
+ integer :: a(:)
+ integer :: b(*)
+ integer, pointer, intent(in) :: c(:)
+
+ !$omp target data map(a,b(:N),c) use_device_addr(a,b(:N),c)
+ !$omp target has_device_addr(A,B(:N),C)
+ if (lbound(a,dim=1) /= 1 .or. ubound(a,dim=1) /= N) stop 10
+ if (lbound(b,dim=1) /= 1) stop 11
+ if (lbound(c,dim=1) /= N .or. ubound(c,dim=1) /= 2*N-1) stop 12
+ if (any (a /= 42)) stop 13
+ if (any (b(:N) /= 43)) stop 14
+ if (any (c /= 44)) stop 15
+ a = [(i, i=1, N)]
+ b(:N) = [(2*i, i = 1, N)]
+ c = [(3*i, i = 1, N)]
+ !$omp end target
+ !$omp end target data
+ end subroutine foo
+
+end program main
diff --git a/libgomp/testsuite/libgomp.fortran/target-has-device-addr-3.f90 b/libgomp/testsuite/libgomp.fortran/target-has-device-addr-3.f90
new file mode 100644
index 00000000000..c6293b4de2e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-has-device-addr-3.f90
@@ -0,0 +1,90 @@
+! Test optional dummy arguments in HAS_DEVICE_ADDR.
+
+program main
+ use omp_lib
+ use iso_c_binding
+ implicit none
+
+ integer, target :: x
+ integer, pointer :: ptr
+ integer, parameter :: N=7
+ real :: y1(N), y2(N)
+ integer, target :: y3(N:2*N-1)
+ integer :: i
+
+ x = 24
+ ptr => x
+ y1 = 42.24
+ y2 = 42.24
+ y3 = 42
+
+ call optional_scalar (is_present=.false.)
+ if (x /= 24) stop 1
+
+ call optional_scalar (x, is_present=.true.)
+ if (x /= 42) stop 2
+
+ call optional_ptr (is_present=.false.)
+ if (x /= 42) stop 3
+ if (ptr /= 42) stop 4
+
+ call optional_ptr (ptr, is_present=.true.)
+ if (x /= 84) stop 5
+ if (ptr /= 84) stop 6
+
+ call optional_array (is_present=.false.)
+ if (any (y1 /= [(42.24, i=1, N)])) stop 7
+ if (any (y2 /= [(42.24, i=1, N)])) stop 8
+ if (any (y3 /= [(42, i=1, N)])) stop 9
+
+ call optional_array (y1, y2, y3, is_present=.true.)
+ if (any (y1 /= [(42.24+i, i=1, N)])) stop 10
+ if (any (y2 /= [(42.24+2*i, i=1, N)])) stop 11
+ if (any (y3 /= [(42+3*i, i=1, N)])) stop 12
+
+contains
+ subroutine optional_scalar (a, is_present)
+ integer, optional :: a
+ logical, value :: is_present
+
+ !$omp target data map(a) use_device_addr(a)
+ !$omp target has_device_addr(a)
+ if (is_present) a = 42
+ !$omp end target
+ !$omp end target data
+ end subroutine optional_scalar
+
+ subroutine optional_ptr (a, is_present)
+ integer, pointer, optional :: a
+ logical, value :: is_present
+ !$omp target data map(a) use_device_addr(a)
+ !$omp target has_device_addr(a)
+ if (is_present) a = 84
+ !$omp end target
+ !$omp end target data
+ end subroutine optional_ptr
+
+ subroutine optional_array (a, b, c, is_present)
+ real, optional :: a(:), b(*)
+ integer, optional, pointer, intent(in) :: c(:)
+ logical, value :: is_present
+ integer :: i
+
+ !$omp target data map(a, b(:N), c) use_device_addr(a, b, c)
+ !$omp target has_device_addr(a, b, c)
+ if (is_present) then
+ if (lbound(a,dim=1) /= 1 .or. ubound(a,dim=1) /= N) stop 21
+ if (lbound(b,dim=1) /= 1) stop 22
+ if (lbound(c,dim=1) /= N .or. ubound(c,dim=1) /= 2*N-1) stop 23
+ if (any (a /= [(42.24, i = 1, N)])) stop 24
+ if (any (b(:N) /= [(42.24, i = 1, N)])) stop 25
+ if (any (c /= [(42, i = 1, N)])) stop 26
+ a = [(42.24+i, i=1, N)]
+ b(:N) = [(42.24+2*i, i=1, N)]
+ c = [(42+3*i, i=1, N)]
+ end if
+ !$omp end target
+ !$omp end target data
+ end subroutine optional_array
+
+end program main
diff --git a/libgomp/testsuite/libgomp.fortran/target-has-device-addr-4.f90 b/libgomp/testsuite/libgomp.fortran/target-has-device-addr-4.f90
new file mode 100644
index 00000000000..59d3e3d31dd
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-has-device-addr-4.f90
@@ -0,0 +1,71 @@
+! Test allocatables in HAS_DEVICE_ADDR.
+
+program main
+ use omp_lib
+ use iso_c_binding
+ implicit none
+
+ integer, parameter :: N = 5
+ integer, allocatable :: x
+ integer, allocatable :: y(:)
+ call scalar_dummy (x)
+ call array_dummy (y)
+ call array_dummy_optional (y)
+ call array_dummy_optional ()
+
+contains
+ subroutine scalar_dummy (a)
+ integer, allocatable :: a
+
+ allocate (a)
+ a = 24
+
+ !$omp target data map(a) use_device_addr(a)
+ !$omp target has_device_addr(a)
+ a = 42
+ !$omp end target
+ !$omp end target data
+ if (a /= 42) stop 1
+
+ deallocate (a)
+ end subroutine scalar_dummy
+
+ subroutine array_dummy (a)
+ integer, allocatable :: a(:)
+ integer :: i
+
+ allocate (a(N))
+ a = 42
+
+ !$omp target data map(a) use_device_addr(a)
+ !$omp target has_device_addr(a)
+ a = [(i, i=1, N)]
+ !$omp end target
+ !$omp end target data
+ if (any (a /= [(i, i=1, N)])) stop 2
+
+ deallocate (a)
+ end subroutine array_dummy
+
+ subroutine array_dummy_optional (a)
+ integer, optional, allocatable :: a(:)
+ integer :: i
+
+ if (present (a)) then
+ allocate (a(N))
+ a = 42
+ end if
+
+ !$omp target data map(a) use_device_addr(a)
+ !$omp target has_device_addr(a)
+ if (present (a)) a = [(i, i=1, N)]
+ !$omp end target
+ !$omp end target data
+
+ if (present (a)) then
+ if (any (a /= [(i, i=1, N)])) stop 2
+ deallocate (a)
+ end if
+ end subroutine array_dummy_optional
+
+end program main
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2022-02-10 7:50 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2022-02-10 7:50 [gcc r12-7161] C, C++, Fortran, OpenMP: Add 'has_device_addr' clause to 'target' construct Marcel Vollweiler
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).