Re: [PATCH, OpenACC 2.7] readonly modifier support in front-ends

[Tobias Burnus <tobias@codesourcery.com>]

Hi Thomas & Chung-Lin,


On 20.07.23 15:33, Thomas Schwinge wrote:
> On 2023-07-11T02:33:58+0800, Chung-Lin Tang
> <chunglin.tang@siemens.com> wrote:
>> +++ b/gcc/c/c-parser.cc
>> @@ -14059,7 +14059,8 @@ c_parser_omp_variable_list (c_parser *parser,
>>
>>   static tree
>>   c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
>> -                           tree list, bool allow_deref = false)
>> +                           tree list, bool allow_deref = false,
>> +                           bool *readonly = NULL)
>> ...
> Instead of doing this in 'c_parser_omp_var_list_parens', I think it's
> clearer to have this special 'readonly :' parsing logic in the two places
> where it's used.

I concur. The same issue also occurred for OpenMP's
c_parser_omp_clause_to, and c_parser_omp_clause_from and the 'present'
modifier. For it, I created a combined function but the main reason for
that is that OpenMP also permits more modifiers (like 'iterators'),
which would cause more duplication of code ('iterator' is not yet
supported).

For something as simple to parse as this modifier, I would just do it at
the two places – as Thomas suggested.

>> +++ b/gcc/fortran/gfortran.h
>> @@ -1360,7 +1360,11 @@ typedef struct gfc_omp_namelist
>>       {
>>         gfc_omp_reduction_op reduction_op;
>>         gfc_omp_depend_doacross_op depend_doacross_op;
>> -      gfc_omp_map_op map_op;
>> +      struct
>> +        {
>> +       ENUM_BITFIELD (gfc_omp_map_op) map_op:8;
>> +       bool readonly;
>> +        };
>>         gfc_expr *align;
>>         struct
>>        {
> [...] Thus, the above looks good to me.

I concur but I wonder whether it would be cleaner to name the struct;
this makes it also more obvious what belongs together in the union.

Namely, naming the struct 'map' and then changing the 45 users from
'u.map_op' to 'u.map.op' and the new 'u.readonly' to 'u.map.readonly'. –
this seems to be cleaner.

>>   static bool
>>   gfc_match_omp_map_clause (gfc_omp_namelist **list, gfc_omp_map_op map_op,
>> -                       bool allow_common, bool allow_derived)
>> +                       bool allow_common, bool allow_derived, bool readonly = false)
>>   {
> Similar to 'c_parser_omp_var_list_parens' above,

I concur that not doing it here is cleaner.

> again, for
> example (random), like 'ancestor :', or 'conditional :' are parsed --
> which you're mostly already doing

I think OpenMP's "present" (as modifier to "omp target updates"'s
"to"/"from") is a better example than "ancestor" as for present we also
have a list. See: gfc_match_motion_var_list how to handle the headp.

(There an extra functions was used as in the future also other modifiers
like 'iterator' will be used.)

However, as Thomas noted, the patch contains also an example (see
further down in Thomas' email, not quoted here).

> Or, we could add a new 'gcc/fortran/gfortran.h:gfc_omp_map_op' item
> 'OMP_MAP_TO_READONLY', which eventually translates into 'OMP_MAP_TO' with
> 'readonly' set?

I think having the additional flag is easier to understand - and at least
memory wise we do not save memory as it is in a union. The advantage
of not having a union is that accessing the int-enum is faster than accessing
an char-wide bitset enum.

In terms of code changes (and without having a closer look), the two
approaches seems to be be similar.

Hence, using OMP_MAP_TO_READONLY for OpenACC would be fine, too. And
I do not have a strong preference for either.

* * *

I did wonder about the following, but I now believe it won't affect
the choice. Namely, we want to handle at some point the following:

!$omp target firstprivate(var) allocator(omp_const_mem_alloc: var)

This could be turned into  GOMP_MAP_FIRSTPRIVATE... + OMP_.*READONLY flag.

But if we don't do it in the FE, the internal Fortran representation
does not matter.
Advantage for doing it in the ME: Only one code location, especially as
we might use the opportunity to also check that the omp_const_mem_alloc
is only used with privatization (in OpenMP).

Difference: OpenMP uses 'firstprivate' (i.e. private copy, no reference count bump,
only permitted for 'target') while OpenACC uses 'copy' which implies reference
counting and permitted in 'acc (enter/exit) data' and not only for compute constructs.

OpenMP in principle also permits user-defined allocator with a constant
memory space - I am not completely sure whether/when it can be used with
   omp target firstprivate(...) allocator(my_alloc : ...)


> Then we'd just here call the (unaltered)
> 'gfc_match_omp_map_clause', with
> 'readonly ? OMP_MAP_TO_READONLY : OMP_MAP_TO'?  Per
> 'git grep --cached '[^G]OMP_MAP_TO[^F]' -- gcc/fortran/' not a lot of
> places need adjusting for that (most of the 'gcc/fortran/openmp.cc' ones
> are not applicable).

I think either would work. – I have no strong feeling what's better.
But you still need to handle it for clause resolution.

> + if (gfc_match ("readonly :") == MATCH_YES)
> I note this one does not have a space after ':' in 'gfc_match', but the
> one above in 'gfc_match_omp_clauses' does.  I don't know off-hand if that
> makes a difference in parsing -- probably not, as all of
> 'gcc/fortran/openmp.cc' generally doesn't seem to be very consistent
> about these two variants?

It *does* make a difference. And for obvious reasons. You don't want to permit:

   !$acc kernels asnyccopy(a)

but require at least one space (or comma) between "async" and "copy"..
(In fixed form Fortran, it would be fine - as would be "!$acc k e nelsasy nc co p y(a)".)

A " " matches zero or more whitespaces, but with gfc_match_space you can find out
whether there was whitespace or not.

Whether the tailing " " in the gfc_match matters or not, depends on what comes next.
If there is a "gfc_gobble_whitespace ();", everything is fine. If not, the next to match
has to start with a " ", which is usually ugly; an exception is " , " or " )" which still
is somewhat fine.

I think that it is mostly implemented correctly, but I wouldn't be surprised if a
space is missing in some matches - be it a tailing white space or e.g. in "foo:" before
the colon.

BTW: One reason of stripping tailing spaces before matching a non-whitespace: the
associated location is the one before the parsing; thus, for a match error or when saving
the old_locus, pointing to the first non-whitespace looks nicer than pointing to the
(first of the) whitspace character(s).


>> +  if (readonly)
>> +    for (gfc_omp_namelist *n = *head; n; n = n->next)
>> +      n->u.readonly = true;
> This already looks like how I thought it should look like.
Indeed.--- a/gcc/tree.h
>> +++ b/gcc/tree.h
>> @@ -1813,6 +1813,14 @@ class auto_suppress_location_wrappers
>>   #define OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE(NODE) \
>>     (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.addressable_flag)
>>
>> +/* Nonzero if OpenACC 'readonly' modifier set, used for 'copyin'.  */
>> +#define OMP_CLAUSE_MAP_READONLY(NODE) \
>> +  TREE_READONLY (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
>> +
>> +/* Same as above, for use in OpenACC cache directives.  */
>> +#define OMP_CLAUSE__CACHE__READONLY(NODE) \
>> +  TREE_READONLY (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE__CACHE_))
> I'm not sure if these special accessor functions are actually useful, or
> we should just directly use 'TREE_READONLY' instead?  We're only using
> them in contexts where it's clear that the 'OMP_CLAUSE_SUBCODE_CHECK' is
> satisfied, for example.
I find directly using TREE_READONLY confusing.
> Setting 'TREE_READONLY' of the 'OMP_CLAUSE_DECL' instead of the clause
> itself isn't the right thing to do -- or is it, and might already
> indicate to the middle end the desired semantics?  But does it maybe
> conflict with front end/language-level use of 'TREE_READONLY' for 'const'
> etc. (I suppose), and thus diagnostics for mismatches?

I think is is cleaner not to one flag to mean two different things.

In particular, wouldn't the following cause issues, if you mark 'a' as TREE_READONLY?

int a;
#pragma acc parallel copyin(readonly : a)
{...}
a = 5;

> Or, early in the middle end, propagate 'TREE_READONLY' from the clause to
> its 'OMP_CLAUSE_DECL'?  Might need to 'unshare_expr' the latter for
> modification and use in the associated region only?

Unsharing a tree would surely help – but it is still ugly and, for
declarations, unshare_expr does not create a copy!

Tobias

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