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

[Thomas Schwinge <thomas@codesourcery.com>]

Hi!

On 2023-08-07T21:58:27+0800, Chung-Lin Tang <chunglin.tang@siemens.com> wrote:
> here's the updated v2 of the readonly modifier front-end patch.

Thanks.


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

> On 2023/7/20 11:08 PM, Tobias Burnus wrote:
>> I concur. [...]
>
> Okay, I've changed the C/C++ parser parts to have the parsing logic directly
> added.

These parts now looks good to me, with one remark for the C front end
changes, see below.


>>>> +++ 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.
>
> I've adjusted 'u.map' to be a named struct now, and updated the references.

I like that, thanks.  (Tobias, to reduce the volume of this patch here,
please let us know if the 'map_op' -> 'map.op' mass-change should be done
separately and go into master branch already, instead of as part of this
patch.)


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

OK, I generally follow -- but does this rationale also apply in this case
here, concerning space after ':'?

> Okay, made sure both are 'gfc_match ("readonly : ")'. Thanks for catching that, didn't
> realize that space was significant.


>>>> +++ 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.
>
> FWIW, I've changed to use TREE_NOTHROW instead, if it can give a better sense of safety :P

I don't understand that, why not use 'TREE_READONLY'?

> I think there's a misunderstanding here anyways: we are not relying on a DECL marked
> TREE_READONLY here. We merely need the OMP_CLAUSE_MAP to be marked as OMP_CLAUSE_MAP_READONLY == 1.

Yes, I understand that.  My question was why we don't just use
'TREE_READONLY (c)', where 'c' is the
'OMP_CLAUSE_MAP'/'OMP_CLAUSE__CACHE_' clause (not its decl), and avoid
the indirection through
'#define OMP_CLAUSE_MAP_READONLY'/'#define OMP_CLAUSE__CACHE__READONLY',
given that we're only using them in contexts where it's clear that the
'OMP_CLAUSE_SUBCODE_CHECK' is satisfied.  I don't have a strong
preference, though.

Either way, you still need to document this:

| Also, for the new use for OMP clauses, update 'gcc/tree.h:TREE_READONLY',
| and in 'gcc/tree-core.h' for 'readonly_flag' the
| "table lists the uses of each of the above flags".


Then, my idea of "Setting 'TREE_READONLY' of the 'OMP_CLAUSE_DECL'
instead of the clause itself" was just that: an idea, so if you conclude
that doesn't make sense, don't follow it further.  In particular, Tobias
said:

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

Aha, my thinking was that we'd have a separate decl inside the compute
region, that is, the host-side 'a' not affected by the 'readonly'
modifier, and thus host-side 'a = 5;' continue to work as expected.

But you're of course right: we cannot set 'TREE_READONLY' early (front
end, before OMP function split off), for the very reason you've cited.
So we definitely need a separate flag, and then it's probably easier
(less invasive) to have it on the clause instead of its decl.  (... as
you've implemented.)

As I said:

| Just some quick thoughts, obviously without any detailed analysis.  ;-)


Another thing, I did wonder: there are cases where for one source-level
OpenACC clause we synthesize several actual clauses (in the front ends,
but possibly also during gimplification?).  Do we understand how such
additionally synthesized clause react to an original clause's 'readonly'
modifier (that is, do they get it propagated, do they also get
'OMP_CLAUSE_MAP_READONLY'/'OMP_CLAUSE__CACHE__READONLY' set, or not?),
and test cases to verify/document that?

Later I found that's part of your follow-on
"[PATCH, OpenACC 2.7] readonly modifier support in front-ends", as you've
also written here:

> The other points-to patch then (also in front-ends) take the OMP_CLAUSE_MAP_READONLY
> to mark the clauses of "base-pointers of array-sections" as OMP_CLAUSE_MAP_POINTS_TO_READONLY,
> and later this gradually gets relayed to alias oracle routines in tree-ssa-alias.cc


> Re-tested this v2 patch on powerpc64le-linux/nvptx. Okay for trunk?

In addition to a few individual comments above and below, you've also not
yet responded to my requests re test cases.


> --- a/gcc/c/c-parser.cc
> +++ b/gcc/c/c-parser.cc
> @@ -14084,7 +14084,11 @@ c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
>     OpenACC 2.6:
>     no_create ( variable-list )
>     attach ( variable-list )
> -   detach ( variable-list ) */
> +   detach ( variable-list )
> +
> +   OpenACC 2.7:
> +   copyin (readonly : variable-list )
> + */
>
>  static tree
>  c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
> @@ -14135,11 +14139,36 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
>      default:
>        gcc_unreachable ();
>      }
> -  tree nl, c;
> -  nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list, true);
>
> -  for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
> -    OMP_CLAUSE_SET_MAP_KIND (c, kind);
> +  tree nl = list;
> +  bool readonly = false;
> +  matching_parens parens;
> +  if (parens.require_open (parser))
> +    {
> +      /* Turn on readonly modifier parsing for copyin clause.  */
> +      if (c_kind == PRAGMA_OACC_CLAUSE_COPYIN)
> +     {
> +       c_token *token = c_parser_peek_token (parser);
> +       if (token->type == CPP_NAME
> +           && !strcmp (IDENTIFIER_POINTER (token->value), "readonly")
> +           && c_parser_peek_2nd_token (parser)->type == CPP_COLON)
> +         {
> +           c_parser_consume_token (parser);
> +           c_parser_consume_token (parser);
> +           readonly = true;
> +         }
> +     }
> +      location_t loc = c_parser_peek_token (parser)->location;

I suppose 'loc' here now points to after the opening '(' or after the
'readonly :'?  This is different from what 'c_parser_omp_var_list_parens'
does, and indeed, 'c_parser_omp_variable_list' states that "CLAUSE_LOC is
the location of the clause", not the location of the variable-list?  As
this, I suppose, may change diagnostics, please restore the original
behavior.  (This appears to be different in the C++ front end, huh.)

> +      nl = c_parser_omp_variable_list (parser, loc, OMP_CLAUSE_MAP, list, true);
> +      parens.skip_until_found_close (parser);
> +    }
> +
> +  for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
> +    {
> +      OMP_CLAUSE_SET_MAP_KIND (c, kind);
> +      if (readonly)
> +     OMP_CLAUSE_MAP_READONLY (c) = 1;
> +    }
>
>    return nl;
>  }
> @@ -18161,15 +18190,40 @@ c_parser_omp_structured_block (c_parser *parser, bool *if_p)
>  /* OpenACC 2.0:
>     # pragma acc cache (variable-list) new-line
>
> +   OpenACC 2.7:
> +   # pragma acc cache (readonly: variable-list) new-line
> +
>     LOC is the location of the #pragma token.
>  */
>
>  static tree
>  c_parser_oacc_cache (location_t loc, c_parser *parser)
>  {
> -  tree stmt, clauses;
> +  tree stmt, clauses = NULL_TREE;
> +  bool readonly = false;
> +  matching_parens parens;
> +
> +  if (parens.require_open (parser))
> +    {
> +      c_token *token = c_parser_peek_token (parser);
> +      if (token->type == CPP_NAME
> +       && !strcmp (IDENTIFIER_POINTER (token->value), "readonly")
> +       && c_parser_peek_2nd_token (parser)->type == CPP_COLON)
> +     {
> +       c_parser_consume_token (parser);
> +       c_parser_consume_token (parser);
> +       readonly = true;
> +     }
> +      location_t loc = c_parser_peek_token (parser)->location;

Similar.  (That is, here, location of the directive.)

> +      clauses = c_parser_omp_variable_list (parser, loc, OMP_CLAUSE__CACHE_,
> +                                         NULL_TREE);
> +      parens.skip_until_found_close (parser);
> +    }
> +
> +  if (readonly)
> +    for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
> +      OMP_CLAUSE__CACHE__READONLY (c) = 1;
>
> -  clauses = c_parser_omp_var_list_parens (parser, OMP_CLAUSE__CACHE_, NULL);
>    clauses = c_finish_omp_clauses (clauses, C_ORT_ACC);
>
>    c_parser_skip_to_pragma_eol (parser);


> --- a/gcc/fortran/openmp.cc
> +++ b/gcc/fortran/openmp.cc
> @@ -1197,7 +1197,7 @@ omp_inv_mask::omp_inv_mask (const omp_mask &m) : omp_mask (m)
>
>  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)
>  {
>    gfc_omp_namelist **head = NULL;
>    if (gfc_match_omp_variable_list ("", list, allow_common, NULL, &head, true,
> @@ -1206,7 +1206,10 @@ gfc_match_omp_map_clause (gfc_omp_namelist **list, gfc_omp_map_op map_op,
>      {
>        gfc_omp_namelist *n;
>        for (n = *head; n; n = n->next)
> -     n->u.map_op = map_op;
> +     {
> +       n->u.map.op = map_op;
> +       n->u.map.readonly = readonly;
> +     }
>        return true;
>      }

Didn't we conclude that "not doing it here is cleaner" (Tobias' words),
and instead do this "Similar to 'c_parser_omp_var_list_parens'" (my
words)?  That is, not add the 'bool readonly' formal parameter to
'gfc_match_omp_map_clause'.

(..., but don't do the 'OMP_MAP_TO_READONLY' way that I considered, but
instead keep the 'readonly' flag.)


Grüße
 Thomas
-----------------
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