Hi Jakub,
Am 20.10.2021 um 14:38 schrieb Jakub Jelinek:
On Mon, Oct 18, 2021 at 06:17:20PM +0200, Marcel Vollweiler wrote:
@@ -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)
OpenMP 5.1 in [200:6-9] says:
The has_device_addr clause indicates ... The list items may include array
sections.
This means in addition to the c-parser.c and parser.c changes you've done,
at least c_parser_omp_variable_list needs to change to include
OMP_CLAUSE_HAS_DEVICE_ADDR among
case OMP_CLAUSE_AFFINITY:
case OMP_CLAUSE_DEPEND:
case OMP_CLAUSE_REDUCTION:
case OMP_CLAUSE_IN_REDUCTION:
case OMP_CLAUSE_TASK_REDUCTION:
clauses (similarly for C++) and then {,c_}finish_omp_clauses needs to handle
it similarly to other clauses that can have array sections.
As it is a data sharing clause, I think the closest model (e.g. for
handle_omp_array_sections* purposes) is OMP_CLAUSE_*REDUCTION.
Then even the case when OMP_CLAUSE_DECL of the clause needs handling
similarly to other clauses that accept array sections.
The handling for array sections is added now. The basic idea of the
implementation is that it seems to be sufficient to consider the base
variable. I'm not completely sure but I think access to memory which is
not specified in has_device_addr cannot be prevented at all and my
reading of the OpenMP 5.1 specification is that the behavour is
undefined for access to memory that is not specified in has_device_addr.
Thus, limitation of an array to some section does not prevent for using
parts of the array outside the specified array section.
Moreover, cases like
#pragma omp target data map(x[2:3]) use_device_addr(x)
#pragma omp target has_device_addr(x[2:3])
or
#pragma omp target data map(x[2:3]) use_device_addr(x[2:3])
#pragma omp target has_device_addr(x[2:3])
do not work yet, since the use_device_addr clause does currently not
support array sections.
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;
I'd prefer not to add new bitmaps unless necessary, can't the clause use the
same bitmap together with is_device_ptr clause? One can't specify something
both as is_device_ptr and has_device_addr at the same time...
Both bitmaps are now combined to one. I previously seperated the bitmaps
in order to have a clearer naming. Now I called it 'is_on_device' to be
more general than with is_device_ptr or has_device_addr. However, other
suggestions are welcome :)
--- 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,
For C++, another thing is whether the clause should accept also non-static
data members in non-static member functions.
5.1 had in List Item Privatization
"A variable that is part of another variable (as an array or structure element)
cannot be privatized
except if the data-sharing attribute clause is associated with a construct
within a class non-static
member function and the variable is an accessible data member of the object for
which the
non-static member function is invoked."
but I believe that hopefully that can't be applied to has_device_addr which
wasn't declared as data sharing clause (and it really is not in the sense
that it doesn't privatize anything).
But 5.2 moves that stuff to general spot where it applies to all clauses:
"Unless otherwise specified, a variable that is part of another variable (as an
array element or a
structure element) cannot be a variable list item, an extended list item or
locator list item
except if the list appears on a clause that is associated with a construct
within a class
non-static member function and the variable is an accessible data member of the
object for
which the non-static member function is invoked."
has_device_addr has the restriction that the list item already has to have
device address, so I bet the whole class object would need to appear on the
device already, but still it is unclear to me what it would mean there.
Let's ignore that for now.
So, for additional testsuite coverage, e.g.
--- /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);
Allocate 30*sizeof(int) and arrange
+
+ #pragma omp target is_device_ptr(dp)
+ *dp = 42;
+
+ int &x = *dp;
For x to be int (&x)[30];
and then test all of has_device_addr(x), has_device_addr(x[3]),
has_device_addr(x[0:22]), has_device_addr(x[17:]) etc.
Similarly for cases with use_device_addr.
Some test cases were added. I'm not sure if array sections should work
on reference types? From the code point of view,
'handle_omp_array_sections_1' rejects explicitly everything which is not
of type array or pointer.
Jakub
Thanks,
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
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_variable_list): Handle array sections.
(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 (handle_omp_array_sections): Handle clause restrictions.
(c_finish_omp_clauses): Handle array sections.
gcc/cp/ChangeLog:
* parser.c (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): 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 (handle_omp_array_sections): Handle clause restrictions.
(finish_omp_clauses): Handle array sections.
gcc/ChangeLog:
* gimplify.c (gimplify_scan_omp_clauses): Add
OMP_CLAUSE_HAS_DEVICE_ADDR cases
and handle array sections.
(gimplify_adjust_omp_clauses): Add OMP_CLAUSE_HAS_DEVICE_ADDR case.
* omp-low.c (scan_sharing_clauses): Handle OMP_CLAUSE_HAS_DEVICE_ADDR.
(lower_omp_target): Same.
* tree-core.h (enum omp_clause_code): Same.
* tree-nested.c (convert_nonlocal_omp_clauses): Same.
(convert_local_omp_clauses): Same.
* tree-pretty-print.c (dump_omp_clause): Same.
* tree.c: Same.
libgomp/ChangeLog:
* libgomp.texi:
* 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-c++-common/target-has-device-addr-1.c: New test.
* testsuite/libgomp.c/target-has-device-addr-3.c: 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-is-device-ptr.c: New test.
diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c
index fad0606..2f38b38 100644
--- a/gcc/c-family/c-omp.c
+++ b/gcc/c-family/c-omp.c
@@ -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 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 80dd61d..a70302d 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;
@@ -13128,6 +13130,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:
while (c_parser_next_token_is (parser, CPP_OPEN_SQUARE))
{
tree low_bound = NULL_TREE, length = NULL_TREE;
@@ -14255,6 +14258,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 +16958,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 +20943,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 782414f..4f238a2 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -13765,6 +13765,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;
@@ -14071,7 +14073,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;
@@ -14106,6 +14108,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))
@@ -14534,7 +14537,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)
@@ -15099,7 +15104,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),
@@ -15184,6 +15190,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
@@ -15204,6 +15212,23 @@ 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);
+ }
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
+ 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.c b/gcc/cp/parser.c
index 32de97b..17da0b7 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -36159,7 +36159,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;
@@ -36450,6 +36452,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:
while (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_SQUARE))
{
tree low_bound = NULL_TREE, length = NULL_TREE;
@@ -39844,6 +39847,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);
@@ -44019,7 +44027,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 b2916f8..50ddd93 100644
--- a/gcc/cp/pt.c
+++ b/gcc/cp/pt.c
@@ -17431,6 +17431,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)
@@ -17570,6 +17571,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 2443d03..be9303b 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -5602,6 +5602,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 (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
|| (TREE_CODE (t) == COMPONENT_REF
@@ -6607,7 +6609,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;
@@ -6639,6 +6641,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))
@@ -6937,7 +6940,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)
@@ -8059,7 +8064,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),
@@ -8313,6 +8319,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);
@@ -8342,6 +8350,23 @@ 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);
+ }
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
+ 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/gimplify.c b/gcc/gimplify.c
index c2ab96e..107e272 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -10031,6 +10031,15 @@ 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);
+ if (TREE_CODE (decl) == ARRAY_REF)
+ {
+ flags = GOVD_FIRSTPRIVATE | GOVD_EXPLICIT;
+ while (TREE_CODE (decl) == ARRAY_REF)
+ decl = TREE_OPERAND (decl, 0);
+ goto do_add_decl;
+ }
case OMP_CLAUSE_IS_DEVICE_PTR:
flags = GOVD_FIRSTPRIVATE | GOVD_EXPLICIT;
goto do_add;
@@ -11459,6 +11468,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 f58a191..ad25ed9 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -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,7 +1384,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);
@@ -1452,6 +1454,11 @@ 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);
+ decl = get_base_address (decl);
+ goto do_private;
+
case OMP_CLAUSE_IS_DEVICE_PTR:
decl = OMP_CLAUSE_DECL (c);
goto do_private;
@@ -1729,12 +1736,17 @@ 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)
+ decl = get_base_address (decl);
+
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);
@@ -12813,8 +12825,11 @@ 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)
+ var = get_base_address (var);
map_cnt++;
if (is_variable_sized (var))
{
@@ -12829,7 +12844,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))
@@ -13282,17 +13298,22 @@ 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)
+ ovar = get_base_address (ovar);
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);
@@ -13314,7 +13335,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)
@@ -13329,6 +13351,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))))
@@ -13526,6 +13549,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;
@@ -13536,12 +13560,17 @@ 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)
+ var = get_base_address (var);
+ x = build_receiver_ref (var, false, ctx);
+ }
if (is_array_data)
{
@@ -13588,7 +13617,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..be0c8db
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-has-device-addr-1.c
@@ -0,0 +1,59 @@
+/* { 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)
+ ;
+
+}
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 f0c65a2..9253ba5 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -345,6 +345,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.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..e7fb1be 100644
--- a/gcc/tree.c
+++ b/gcc/tree.c
@@ -306,6 +306,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 */
@@ -395,6 +396,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 fd747b9..c77d5e9 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..d9a309d7
--- /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 0000000..401a30d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-has-device-addr-4.C
@@ -0,0 +1,24 @@
+#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 ();
+
+ 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..12040dc
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-has-device-addr-1.c
@@ -0,0 +1,65 @@
+/* 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)
+ ;
+
+ 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 0000000..fd99a82
--- /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;
+}
