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.
> 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...
> --- 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.
Jakub
