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

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

[Thomas Schwinge]

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

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

[Chung-Lin Tang]

[-- Attachment #1: Type: text/plain, Size: 7876 bytes --]

Hi Thomas, Tobias,

On 2023/10/26 6:43 PM, Thomas Schwinge wrote:
>>>>> +++ 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.

After further re-testing using TREE_NOTHROW, I have reverted to using TREE_READONLY, because TREE_NOTHROW clashes
with OMP_CLAUSE_RELEASE_DESCRIPTOR (which doesn't use the OMP_CLAUSE_MAP_* naming convention and is
not documented in gcc/tree-core.h either, hmmm...)

I have added the comment adjustments in gcc/tree-core.h for the new uses of TREE_READONLY/readonly_flag.

We basically all use OMP_CLAUSE_SUBCODE_CHECK macros for OpenMP clause expressions exclusively,
so I don't see a reason to diverge from that style (even when context is clear).

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

Okay, done as mentioned above.

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

I have greatly expanded the test scan patterns to include parallel/kernels/serial/data/enter data,
as well as non-readonly copyin clause together with readonly.

Also added simple 'declare' tests, but there is not anything to scan in the 'tree-original' dump though.

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

Thanks for catching this! Fixed.

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

Fixed in this v3 patch.

Again, tested on x86_64-linux + nvptx offloading. Okay for mainline?

Thanks,
Chung-Lin

gcc/c/ChangeLog:

	* c-parser.cc (c_parser_oacc_data_clause): Add parsing support for
	'readonly' modifier, set OMP_CLAUSE_MAP_READONLY if readonly modifier
	found, update comments.
	(c_parser_oacc_cache): Add parsing support for 'readonly' modifier,
	set OMP_CLAUSE__CACHE__READONLY if readonly modifier found, update
	comments.

gcc/cp/ChangeLog:

	* parser.cc (cp_parser_oacc_data_clause): Add parsing support for
	'readonly' modifier, set OMP_CLAUSE_MAP_READONLY if readonly modifier
	found, update comments.
	(cp_parser_oacc_cache): Add parsing support for 'readonly' modifier,
	set OMP_CLAUSE__CACHE__READONLY if readonly modifier found, update
	comments.

gcc/fortran/ChangeLog:

	* dump-parse-tree.cc (show_omp_namelist): Print "readonly," for
	OMP_LIST_MAP and OMP_LIST_CACHE if n->u.map.readonly is set.
	Adjust 'n->u.map_op' to 'n->u.map.op'.
	* gfortran.h (typedef struct gfc_omp_namelist): Adjust map_op as
	'ENUM_BITFIELD (gfc_omp_map_op) op:8', add 'bool readonly' field,
	change to named struct field 'map'.

	* openmp.cc (gfc_match_omp_map_clause): Adjust 'n->u.map_op' to
	'n->u.map.op'.
	(gfc_match_omp_clause_reduction): Likewise.

	(gfc_match_omp_clauses): Add readonly modifier parsing for OpenACC
	copyin clause, set 'n->u.map.op' and 'n->u.map.readonly' for parsed
	clause. Adjust 'n->u.map_op' to 'n->u.map.op'.
	(gfc_match_oacc_declare): Adjust 'n->u.map_op' to 'n->u.map.op'.
	(gfc_match_oacc_cache): Add readonly modifier parsing for OpenACC
	cache directive.
	(resolve_omp_clauses): Adjust 'n->u.map_op' to 'n->u.map.op'.
	* trans-decl.cc (add_clause): Adjust 'n->u.map_op' to 'n->u.map.op'.
	(finish_oacc_declare): Likewise.
	* trans-openmp.cc (gfc_trans_omp_clauses): Set OMP_CLAUSE_MAP_READONLY,
	OMP_CLAUSE__CACHE__READONLY to 1 when readonly is set. Adjust
	'n->u.map_op' to 'n->u.map.op'.
	(gfc_add_clause_implicitly): Adjust 'n->u.map_op' to 'n->u.map.op'.

gcc/ChangeLog:
	* tree.h (OMP_CLAUSE_MAP_READONLY): New macro.
	(OMP_CLAUSE__CACHE__READONLY): New macro.
	* tree-core.h (struct GTY(()) tree_base): Adjust comments for new
	uses of readonly_flag bit in OMP_CLAUSE_MAP_READONLY and
	OMP_CLAUSE__CACHE__READONLY.
	* tree-pretty-print.cc (dump_omp_clause): Add support for printing
	OMP_CLAUSE_MAP_READONLY and OMP_CLAUSE__CACHE__READONLY.

gcc/testsuite/ChangeLog:

	* c-c++-common/goacc/readonly-1.c: New test.
	* gfortran.dg/goacc/readonly-1.f90: New test.






[-- Attachment #2: readonly-fe-v3.patch --]
[-- Type: text/plain, Size: 32644 bytes --]

diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 53e99aa29d9..00f8bf4376e 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -15627,7 +15627,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,
@@ -15680,11 +15684,37 @@ 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, false);
 
-  for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
-    OMP_CLAUSE_SET_MAP_KIND (c, kind);
+  tree nl = list;
+  bool readonly = false;
+  location_t open_loc = c_parser_peek_token (parser)->location;
+  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;
+	    }
+	}
+      nl = c_parser_omp_variable_list (parser, open_loc, OMP_CLAUSE_MAP, list,
+				       false);
+      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;
 }
@@ -19821,15 +19851,39 @@ 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;
+  location_t open_loc = c_parser_peek_token (parser)->location;
+  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;
+	}
+      clauses = c_parser_omp_variable_list (parser, open_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);
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index e32acfc30a2..4fe27fb07b2 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -38544,7 +38544,11 @@ cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list,
    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
 cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
@@ -38597,11 +38601,34 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
     default:
       gcc_unreachable ();
     }
-  tree nl, c;
-  nl = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list, false);
 
-  for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
-    OMP_CLAUSE_SET_MAP_KIND (c, kind);
+  tree nl = list;
+  bool readonly = false;
+  if (cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
+    {
+      /* Turn on readonly modifier parsing for copyin clause.  */
+      if (c_kind == PRAGMA_OACC_CLAUSE_COPYIN)
+	{
+	  cp_token *token = cp_lexer_peek_token (parser->lexer);
+	  if (token->type == CPP_NAME
+	      && !strcmp (IDENTIFIER_POINTER (token->u.value), "readonly")
+	      && cp_lexer_peek_nth_token (parser->lexer, 2)->type == CPP_COLON)
+	    {
+	      cp_lexer_consume_token (parser->lexer);
+	      cp_lexer_consume_token (parser->lexer);
+	      readonly = true;
+	    }
+	}
+      nl = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE_MAP, list, NULL,
+					   false);
+    }
+
+  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;
 }
@@ -47178,6 +47205,9 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
 
 /* OpenACC 2.0:
    # pragma acc cache (variable-list) new-line
+
+   OpenACC 2.7:
+   # pragma acc cache (readonly: variable-list) new-line
 */
 
 static tree
@@ -47187,9 +47217,28 @@ cp_parser_oacc_cache (cp_parser *parser, cp_token *pragma_tok)
      clauses.  */
   auto_suppress_location_wrappers sentinel;
 
-  tree stmt, clauses;
+  tree stmt, clauses = NULL_TREE;
+  bool readonly = false;
+
+  if (cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
+    {
+      cp_token *token = cp_lexer_peek_token (parser->lexer);
+      if (token->type == CPP_NAME
+	  && !strcmp (IDENTIFIER_POINTER (token->u.value), "readonly")
+	  && cp_lexer_peek_nth_token (parser->lexer, 2)->type == CPP_COLON)
+	{
+	  cp_lexer_consume_token (parser->lexer);
+	  cp_lexer_consume_token (parser->lexer);
+	  readonly = true;
+	}
+      clauses = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE__CACHE_,
+						NULL, NULL);
+    }
+
+  if (readonly)
+    for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+      OMP_CLAUSE__CACHE__READONLY (c) = 1;
 
-  clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE__CACHE_, NULL_TREE);
   clauses = finish_omp_clauses (clauses, C_ORT_ACC);
 
   cp_parser_require_pragma_eol (parser, cp_lexer_peek_token (parser->lexer));
diff --git a/gcc/fortran/dump-parse-tree.cc b/gcc/fortran/dump-parse-tree.cc
index 7b154eb3ca7..db84b06289b 100644
--- a/gcc/fortran/dump-parse-tree.cc
+++ b/gcc/fortran/dump-parse-tree.cc
@@ -1400,6 +1400,9 @@ show_omp_namelist (int list_type, gfc_omp_namelist *n)
 	    fputs (") ALLOCATE(", dumpfile);
 	  continue;
 	}
+      if ((list_type == OMP_LIST_MAP || list_type == OMP_LIST_CACHE)
+	  && n->u.map.readonly)
+	fputs ("readonly,", dumpfile);
       if (list_type == OMP_LIST_REDUCTION)
 	switch (n->u.reduction_op)
 	  {
@@ -1467,7 +1470,7 @@ show_omp_namelist (int list_type, gfc_omp_namelist *n)
 	  default: break;
 	  }
       else if (list_type == OMP_LIST_MAP)
-	switch (n->u.map_op)
+	switch (n->u.map.op)
 	  {
 	  case OMP_MAP_ALLOC: fputs ("alloc:", dumpfile); break;
 	  case OMP_MAP_TO: fputs ("to:", dumpfile); break;
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index ebba2336e12..32b792f85fb 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1363,7 +1363,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) op:8;
+	  bool readonly;
+        } map;
       gfc_expr *align;
       struct
 	{
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index 38de60238c0..5c44e666eb9 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -1210,7 +1210,7 @@ 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;
       return true;
     }
 
@@ -1524,7 +1524,7 @@ gfc_match_omp_clause_reduction (char pc, gfc_omp_clauses *c, bool openacc,
 	    gfc_omp_namelist *p = gfc_get_omp_namelist (), **tl;
 	    p->sym = n->sym;
 	    p->where = p->where;
-	    p->u.map_op = OMP_MAP_ALWAYS_TOFROM;
+	    p->u.map.op = OMP_MAP_ALWAYS_TOFROM;
 
 	    tl = &c->lists[OMP_LIST_MAP];
 	    while (*tl)
@@ -2181,11 +2181,25 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 	    {
 	      if (openacc)
 		{
-		  if (gfc_match ("copyin ( ") == MATCH_YES
-		      && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
-						   OMP_MAP_TO, true,
-						   allow_derived))
-		    continue;
+		  if (gfc_match ("copyin ( ") == MATCH_YES)
+		    {
+		      bool readonly = gfc_match ("readonly : ") == MATCH_YES;
+		      head = NULL;
+		      if (gfc_match_omp_variable_list ("",
+						       &c->lists[OMP_LIST_MAP],
+						       true, NULL, &head, true,
+						       allow_derived)
+			  == MATCH_YES)
+			{
+			  gfc_omp_namelist *n;
+			  for (n = *head; n; n = n->next)
+			    {
+			      n->u.map.op = OMP_MAP_TO;
+			      n->u.map.readonly = readonly;
+			    }
+			  continue;
+			}
+		    }
 		}
 	      else if (gfc_match_omp_variable_list ("copyin (",
 						    &c->lists[OMP_LIST_COPYIN],
@@ -3134,7 +3148,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 		{
 		  gfc_omp_namelist *n;
 		  for (n = *head; n; n = n->next)
-		    n->u.map_op = map_op;
+		    n->u.map.op = map_op;
 		  continue;
 		}
 	      gfc_current_locus = old_loc;
@@ -4002,7 +4016,7 @@ gfc_match_oacc_declare (void)
       if (gfc_current_ns->proc_name
 	  && gfc_current_ns->proc_name->attr.flavor == FL_MODULE)
 	{
-	  if (n->u.map_op != OMP_MAP_ALLOC && n->u.map_op != OMP_MAP_TO)
+	  if (n->u.map.op != OMP_MAP_ALLOC && n->u.map.op != OMP_MAP_TO)
 	    {
 	      gfc_error ("Invalid clause in module with !$ACC DECLARE at %L",
 			 &where);
@@ -4036,7 +4050,7 @@ gfc_match_oacc_declare (void)
 	  return MATCH_ERROR;
 	}
 
-      switch (n->u.map_op)
+      switch (n->u.map.op)
 	{
 	  case OMP_MAP_FORCE_ALLOC:
 	  case OMP_MAP_ALLOC:
@@ -4151,21 +4165,36 @@ gfc_match_oacc_wait (void)
 match
 gfc_match_oacc_cache (void)
 {
+  bool readonly = false;
   gfc_omp_clauses *c = gfc_get_omp_clauses ();
   /* The OpenACC cache directive explicitly only allows "array elements or
      subarrays", which we're currently not checking here.  Either check this
      after the call of gfc_match_omp_variable_list, or add something like a
      only_sections variant next to its allow_sections parameter.  */
-  match m = gfc_match_omp_variable_list (" (",
-					 &c->lists[OMP_LIST_CACHE], true,
-					 NULL, NULL, true);
+  match m = gfc_match (" ( ");
   if (m != MATCH_YES)
     {
       gfc_free_omp_clauses(c);
       return m;
     }
 
-  if (gfc_current_state() != COMP_DO 
+  if (gfc_match ("readonly : ") == MATCH_YES)
+    readonly = true;
+
+  gfc_omp_namelist **head = NULL;
+  m = gfc_match_omp_variable_list ("", &c->lists[OMP_LIST_CACHE], true,
+				   NULL, &head, true);
+  if (m != MATCH_YES)
+    {
+      gfc_free_omp_clauses(c);
+      return m;
+    }
+
+  if (readonly)
+    for (gfc_omp_namelist *n = *head; n; n = n->next)
+      n->u.map.readonly = true;
+
+  if (gfc_current_state() != COMP_DO
       && gfc_current_state() != COMP_DO_CONCURRENT)
     {
       gfc_error ("ACC CACHE directive must be inside of loop %C");
@@ -8436,8 +8465,8 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 		  }
 		if (openacc
 		    && list == OMP_LIST_MAP
-		    && (n->u.map_op == OMP_MAP_ATTACH
-			|| n->u.map_op == OMP_MAP_DETACH))
+		    && (n->u.map.op == OMP_MAP_ATTACH
+			|| n->u.map.op == OMP_MAP_DETACH))
 		  {
 		    symbol_attribute attr;
 		    if (n->expr)
@@ -8447,7 +8476,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 		    if (!attr.pointer && !attr.allocatable)
 		      gfc_error ("%qs clause argument must be ALLOCATABLE or "
 				 "a POINTER at %L",
-				 (n->u.map_op == OMP_MAP_ATTACH) ? "attach"
+				 (n->u.map.op == OMP_MAP_ATTACH) ? "attach"
 				 : "detach", &n->where);
 		  }
 		if (lastref
@@ -8518,7 +8547,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 		else if (openacc)
 		  {
 		    if (list == OMP_LIST_MAP
-			&& n->u.map_op == OMP_MAP_FORCE_DEVICEPTR)
+			&& n->u.map.op == OMP_MAP_FORCE_DEVICEPTR)
 		      resolve_oacc_deviceptr_clause (n->sym, n->where, name);
 		    else
 		      resolve_oacc_data_clauses (n->sym, n->where, name);
@@ -8540,7 +8569,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 		    {
 		    case EXEC_OMP_TARGET:
 		    case EXEC_OMP_TARGET_DATA:
-		      switch (n->u.map_op)
+		      switch (n->u.map.op)
 			{
 			case OMP_MAP_TO:
 			case OMP_MAP_ALWAYS_TO:
@@ -8567,7 +8596,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 			}
 		      break;
 		    case EXEC_OMP_TARGET_ENTER_DATA:
-		      switch (n->u.map_op)
+		      switch (n->u.map.op)
 			{
 			case OMP_MAP_TO:
 			case OMP_MAP_ALWAYS_TO:
@@ -8577,16 +8606,16 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 			case OMP_MAP_PRESENT_ALLOC:
 			  break;
 			case OMP_MAP_TOFROM:
-			  n->u.map_op = OMP_MAP_TO;
+			  n->u.map.op = OMP_MAP_TO;
 			  break;
 			case OMP_MAP_ALWAYS_TOFROM:
-			  n->u.map_op = OMP_MAP_ALWAYS_TO;
+			  n->u.map.op = OMP_MAP_ALWAYS_TO;
 			  break;
 			case OMP_MAP_PRESENT_TOFROM:
-			  n->u.map_op = OMP_MAP_PRESENT_TO;
+			  n->u.map.op = OMP_MAP_PRESENT_TO;
 			  break;
 			case OMP_MAP_ALWAYS_PRESENT_TOFROM:
-			  n->u.map_op = OMP_MAP_ALWAYS_PRESENT_TO;
+			  n->u.map.op = OMP_MAP_ALWAYS_PRESENT_TO;
 			  break;
 			default:
 			  gfc_error ("TARGET ENTER DATA with map-type other "
@@ -8596,7 +8625,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 			}
 		      break;
 		    case EXEC_OMP_TARGET_EXIT_DATA:
-		      switch (n->u.map_op)
+		      switch (n->u.map.op)
 			{
 			case OMP_MAP_FROM:
 			case OMP_MAP_ALWAYS_FROM:
@@ -8606,16 +8635,16 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 			case OMP_MAP_DELETE:
 			  break;
 			case OMP_MAP_TOFROM:
-			  n->u.map_op = OMP_MAP_FROM;
+			  n->u.map.op = OMP_MAP_FROM;
 			  break;
 			case OMP_MAP_ALWAYS_TOFROM:
-			  n->u.map_op = OMP_MAP_ALWAYS_FROM;
+			  n->u.map.op = OMP_MAP_ALWAYS_FROM;
 			  break;
 			case OMP_MAP_PRESENT_TOFROM:
-			  n->u.map_op = OMP_MAP_PRESENT_FROM;
+			  n->u.map.op = OMP_MAP_PRESENT_FROM;
 			  break;
 			case OMP_MAP_ALWAYS_PRESENT_TOFROM:
-			  n->u.map_op = OMP_MAP_ALWAYS_PRESENT_FROM;
+			  n->u.map.op = OMP_MAP_ALWAYS_PRESENT_FROM;
 			  break;
 			default:
 			  gfc_error ("TARGET EXIT DATA with map-type other "
diff --git a/gcc/fortran/trans-decl.cc b/gcc/fortran/trans-decl.cc
index 6d463036966..b7dea11461f 100644
--- a/gcc/fortran/trans-decl.cc
+++ b/gcc/fortran/trans-decl.cc
@@ -6744,7 +6744,7 @@ add_clause (gfc_symbol *sym, gfc_omp_map_op map_op)
 
   n = gfc_get_omp_namelist ();
   n->sym = sym;
-  n->u.map_op = map_op;
+  n->u.map.op = map_op;
 
   if (!module_oacc_clauses)
     module_oacc_clauses = gfc_get_omp_clauses ();
@@ -6846,10 +6846,10 @@ finish_oacc_declare (gfc_namespace *ns, gfc_symbol *sym, bool block)
 
   for (n = omp_clauses->lists[OMP_LIST_MAP]; n; n = n->next)
     {
-      switch (n->u.map_op)
+      switch (n->u.map.op)
 	{
 	  case OMP_MAP_DEVICE_RESIDENT:
-	    n->u.map_op = OMP_MAP_FORCE_ALLOC;
+	    n->u.map.op = OMP_MAP_FORCE_ALLOC;
 	    break;
 
 	  default:
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index a2bf15665b3..fa1bfd41380 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -3139,7 +3139,10 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		      || (n->expr && gfc_expr_attr (n->expr).pointer)))
 		always_modifier = true;
 
-	      switch (n->u.map_op)
+	      if (n->u.map.readonly)
+		OMP_CLAUSE_MAP_READONLY (node) = 1;
+
+	      switch (n->u.map.op)
 		{
 		case OMP_MAP_ALLOC:
 		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC);
@@ -3266,8 +3269,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		      && n->sym->attr.omp_declare_target
 		      && (always_modifier || n->sym->attr.pointer)
 		      && op != EXEC_OMP_TARGET_EXIT_DATA
-		      && n->u.map_op != OMP_MAP_DELETE
-		      && n->u.map_op != OMP_MAP_RELEASE)
+		      && n->u.map.op != OMP_MAP_DELETE
+		      && n->u.map.op != OMP_MAP_RELEASE)
 		    {
 		      gcc_assert (n->sym->ts.u.cl->backend_decl);
 		      node5 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
@@ -3333,7 +3336,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			{
 			  enum gomp_map_kind gmk = GOMP_MAP_POINTER;
 			  if (op == EXEC_OMP_TARGET_EXIT_DATA
-			      && n->u.map_op == OMP_MAP_DELETE)
+			      && n->u.map.op == OMP_MAP_DELETE)
 			    gmk = GOMP_MAP_DELETE;
 			  else if (op == EXEC_OMP_TARGET_EXIT_DATA)
 			    gmk = GOMP_MAP_RELEASE;
@@ -3356,7 +3359,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			{
 			  enum gomp_map_kind gmk;
 			  if (op == EXEC_OMP_TARGET_EXIT_DATA
-			      && n->u.map_op == OMP_MAP_DELETE)
+			      && n->u.map.op == OMP_MAP_DELETE)
 			    gmk = GOMP_MAP_DELETE;
 			  else if (op == EXEC_OMP_TARGET_EXIT_DATA)
 			    gmk = GOMP_MAP_RELEASE;
@@ -3388,18 +3391,18 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		      node2 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
 		      OMP_CLAUSE_DECL (node2) = decl;
 		      OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
-		      if (n->u.map_op == OMP_MAP_DELETE)
+		      if (n->u.map.op == OMP_MAP_DELETE)
 			map_kind = GOMP_MAP_DELETE;
 		      else if (op == EXEC_OMP_TARGET_EXIT_DATA
-			       || n->u.map_op == OMP_MAP_RELEASE)
+			       || n->u.map.op == OMP_MAP_RELEASE)
 			map_kind = GOMP_MAP_RELEASE;
 		      else
 			map_kind = GOMP_MAP_TO_PSET;
 		      OMP_CLAUSE_SET_MAP_KIND (node2, map_kind);
 
 		      if (op != EXEC_OMP_TARGET_EXIT_DATA
-			  && n->u.map_op != OMP_MAP_DELETE
-			  && n->u.map_op != OMP_MAP_RELEASE)
+			  && n->u.map.op != OMP_MAP_DELETE
+			  && n->u.map.op != OMP_MAP_RELEASE)
 			{
 			  node3 = build_omp_clause (input_location,
 						    OMP_CLAUSE_MAP);
@@ -3417,7 +3420,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			      = gfc_conv_descriptor_data_get (decl);
 			  OMP_CLAUSE_SIZE (node3) = size_int (0);
 
-			  if (n->u.map_op == OMP_MAP_ATTACH)
+			  if (n->u.map.op == OMP_MAP_ATTACH)
 			    {
 			      /* Standalone attach clauses used with arrays with
 				 descriptors must copy the descriptor to the
@@ -3433,7 +3436,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			      node3 = NULL;
 			      goto finalize_map_clause;
 			    }
-			  else if (n->u.map_op == OMP_MAP_DETACH)
+			  else if (n->u.map.op == OMP_MAP_DETACH)
 			    {
 			      OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_DETACH);
 			      /* Similarly to above, we don't want to unmap PTR
@@ -3626,8 +3629,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			 to perform a single attach/detach operation, of the
 			 pointer itself, not of the pointed-to object.  */
 		      if (openacc
-			  && (n->u.map_op == OMP_MAP_ATTACH
-			      || n->u.map_op == OMP_MAP_DETACH))
+			  && (n->u.map.op == OMP_MAP_ATTACH
+			      || n->u.map.op == OMP_MAP_DETACH))
 			{
 			  OMP_CLAUSE_DECL (node)
 			    = build_fold_addr_expr (OMP_CLAUSE_DECL (node));
@@ -3656,7 +3659,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 					       se.string_length),
 					   TYPE_SIZE_UNIT (tmp));
 			  gomp_map_kind kind;
-			  if (n->u.map_op == OMP_MAP_DELETE)
+			  if (n->u.map.op == OMP_MAP_DELETE)
 			    kind = GOMP_MAP_DELETE;
 			  else if (op == EXEC_OMP_TARGET_EXIT_DATA)
 			    kind = GOMP_MAP_RELEASE;
@@ -3713,8 +3716,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			     to perform a single attach/detach operation, of the
 			     pointer itself, not of the pointed-to object.  */
 			  if (openacc
-			      && (n->u.map_op == OMP_MAP_ATTACH
-				  || n->u.map_op == OMP_MAP_DETACH))
+			      && (n->u.map.op == OMP_MAP_ATTACH
+				  || n->u.map.op == OMP_MAP_DETACH))
 			    {
 			      OMP_CLAUSE_DECL (node)
 				= build_fold_addr_expr (inner);
@@ -3806,8 +3809,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		    {
 		      /* Bare attach and detach clauses don't want any
 			 additional nodes.  */
-		      if ((n->u.map_op == OMP_MAP_ATTACH
-			   || n->u.map_op == OMP_MAP_DETACH)
+		      if ((n->u.map.op == OMP_MAP_ATTACH
+			   || n->u.map.op == OMP_MAP_DETACH)
 			  && (POINTER_TYPE_P (TREE_TYPE (inner))
 			      || GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (inner))))
 			{
@@ -3840,8 +3843,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			    map_kind = ((GOMP_MAP_ALWAYS_P (map_kind)
 					 || gfc_expr_attr (n->expr).pointer)
 					? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO);
-			  else if (n->u.map_op == OMP_MAP_RELEASE
-				   || n->u.map_op == OMP_MAP_DELETE)
+			  else if (n->u.map.op == OMP_MAP_RELEASE
+				   || n->u.map.op == OMP_MAP_DELETE)
 			    ;
 			  else if (op == EXEC_OMP_TARGET_EXIT_DATA
 				   || op == EXEC_OACC_EXIT_DATA)
@@ -4088,6 +4091,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		}
 	      if (n->u.present_modifier)
 		OMP_CLAUSE_MOTION_PRESENT (node) = 1;
+	      if (list == OMP_LIST_CACHE && n->u.map.readonly)
+		OMP_CLAUSE__CACHE__READONLY (node) = 1;
 	      omp_clauses = gfc_trans_add_clause (node, omp_clauses);
 	    }
 	  break;
@@ -6561,7 +6566,7 @@ gfc_add_clause_implicitly (gfc_omp_clauses *clauses_out,
 	  n2->where = n->where;
 	  n2->sym = n->sym;
 	  if (is_target)
-	    n2->u.map_op = OMP_MAP_TOFROM;
+	    n2->u.map.op = OMP_MAP_TOFROM;
 	  if (tail)
 	    {
 	      tail->next = n2;
diff --git a/gcc/testsuite/c-c++-common/goacc/readonly-1.c b/gcc/testsuite/c-c++-common/goacc/readonly-1.c
new file mode 100644
index 00000000000..34fc92c24d5
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/readonly-1.c
@@ -0,0 +1,59 @@
+/* { dg-additional-options "-fdump-tree-original" } */
+
+struct S
+{
+  int *ptr;
+  float f;
+};
+
+int a[32], b[32];
+#pragma acc declare copyin(readonly: a) copyin(b)
+
+int main (void)
+{
+  int x[32], y[32];
+  struct S s = {x, 0};
+
+  #pragma acc parallel copyin(readonly: x[:32], s.ptr[:16]) copyin(y[:32])
+  {
+    #pragma acc cache (readonly: x[:32])
+    #pragma acc cache (y[:32])
+  }
+
+  #pragma acc kernels copyin(readonly: x[:32], s.ptr[:16]) copyin(y[:32])
+  {
+    #pragma acc cache (readonly: x[:32])
+    #pragma acc cache (y[:32])
+  }
+
+  #pragma acc serial copyin(readonly: x[:32], s.ptr[:16]) copyin(y[:32])
+  {
+    #pragma acc cache (readonly: x[:32])
+    #pragma acc cache (y[:32])
+  }
+
+  #pragma acc data copyin(readonly: x[:32], s.ptr[:16]) copyin(y[:32])
+  {
+    #pragma acc cache (readonly: x[:32])
+    #pragma acc cache (y[:32])
+  }
+
+  #pragma acc enter data copyin(readonly: x[:32], s.ptr[:16]) copyin(y[:32])
+
+  return 0;
+}
+
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */
+
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */
+
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(readonly:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\);$" 4 "original" } } */
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\);$" 4 "original" } } */
diff --git a/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90 b/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90
new file mode 100644
index 00000000000..696ebd08321
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90
@@ -0,0 +1,89 @@
+! { dg-additional-options "-fdump-tree-original" }
+
+subroutine foo (a, n)
+  integer :: n, a(:)
+  integer :: i, b(n), c(n)
+  !$acc parallel copyin(readonly: a(:), b(:n)) copyin(c(:))
+  do i = 1,32
+     !$acc cache (readonly: a(:), b(:n))
+     !$acc cache (c(:))
+  enddo
+  !$acc end parallel
+
+  !$acc kernels copyin(readonly: a(:), b(:n)) copyin(c(:))
+  do i = 1,32
+     !$acc cache (readonly: a(:), b(:n))
+     !$acc cache (c(:))
+  enddo
+  !$acc end kernels
+
+  !$acc serial copyin(readonly: a(:), b(:n)) copyin(c(:))
+  do i = 1,32
+     !$acc cache (readonly: a(:), b(:n))
+     !$acc cache (c(:))
+  enddo
+  !$acc end serial
+
+  !$acc data copyin(readonly: a(:), b(:n)) copyin(c(:))
+  do i = 1,32
+     !$acc cache (readonly: a(:), b(:n))
+     !$acc cache (c(:))
+  enddo
+  !$acc end data
+
+  !$acc enter data copyin(readonly: a(:), b(:n)) copyin(c(:))
+
+end subroutine foo
+
+program main
+  integer :: g(32), h(32)
+  integer :: i, n = 32, a(32)
+  integer :: b(32), c(32)
+
+  !$acc declare copyin(readonly: g), copyin(h)
+
+  !$acc parallel copyin(readonly: a(:32), b(:n)) copyin(c(:))
+  do i = 1,32
+     !$acc cache (readonly: a(:), b(:n))
+     !$acc cache (c(:))
+  enddo
+  !$acc end parallel
+
+  !$acc kernels copyin(readonly: a(:), b(:n)) copyin(c(:))
+  do i = 1,32
+     !$acc cache (readonly: a(:), b(:n))
+     !$acc cache (c(:))
+  enddo
+  !$acc end kernels
+
+  !$acc serial copyin(readonly: a(:), b(:n)) copyin(c(:))
+  do i = 1,32
+     !$acc cache (readonly: a(:), b(:n))
+     !$acc cache (c(:))
+  enddo
+  !$acc end serial
+
+  !$acc data copyin(readonly: a(:), b(:n)) copyin(c(:))
+  do i = 1,32
+     !$acc cache (readonly: a(:), b(:n))
+     !$acc cache (c(:))
+  enddo
+  !$acc end data
+
+  !$acc enter data copyin(readonly: a(:), b(:n)) copyin(c(:))
+
+end program main
+
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(readonly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(readonly,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(readonly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(readonly,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(readonly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(readonly,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(readonly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(readonly,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } }
+
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(readonly:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\) \\(readonly:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\);" 8 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\);" 8 "original" } }
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index 8a89462bd7e..d529712306d 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -1344,6 +1344,12 @@ struct GTY(()) tree_base {
        TYPE_READONLY in
            all types
 
+       OMP_CLAUSE_MAP_READONLY in
+           OMP_CLAUSE_MAP
+
+       OMP_CLAUSE__CACHE__READONLY in
+           OMP_CLAUSE__CACHE_
+
    constant_flag:
 
        TREE_CONSTANT in
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index 654f5247e3a..926f7e006a7 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -913,6 +913,8 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
 
     case OMP_CLAUSE_MAP:
       pp_string (pp, "map(");
+      if (OMP_CLAUSE_MAP_READONLY (clause))
+	pp_string (pp, "readonly,");
       switch (OMP_CLAUSE_MAP_KIND (clause))
 	{
 	case GOMP_MAP_ALLOC:
@@ -1095,6 +1097,8 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
 
     case OMP_CLAUSE__CACHE_:
       pp_string (pp, "(");
+      if (OMP_CLAUSE__CACHE__READONLY (clause))
+	pp_string (pp, "readonly:");
       dump_generic_node (pp, OMP_CLAUSE_DECL (clause),
 			 spc, flags, false);
       goto print_clause_size;
diff --git a/gcc/tree.h b/gcc/tree.h
index e1fc6c2221d..b67a37d6522 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1841,6 +1841,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_))
+
 /* True on an OMP_CLAUSE_USE_DEVICE_PTR with an OpenACC 'if_present'
    clause.  */
 #define OMP_CLAUSE_USE_DEVICE_PTR_IF_PRESENT(NODE) \

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

[Thomas Schwinge]

Hi Chung-Lin!

On 2024-03-07T17:02:02+0900, Chung-Lin Tang <cltang@pllab.cs.nthu.edu.tw> wrote:
> On 2023/10/26 6:43 PM, Thomas Schwinge wrote:
>>>>>> +++ 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.
>
> After further re-testing using TREE_NOTHROW, I have reverted to using TREE_READONLY

ACK, thanks.

> because TREE_NOTHROW clashes
> with OMP_CLAUSE_RELEASE_DESCRIPTOR (which doesn't use the OMP_CLAUSE_MAP_* naming convention and is
> not documented in gcc/tree-core.h either, hmmm...)

Yeah, it's a mess...  The same bits of information spread over three
different places.

(One day I'll turn 'tree's into a proper C++ class hierarchy, with
accessor methods for such flags, statically checked at compile-time, and
thus documented in a single place.  Etc.)

> I have added the comment adjustments in gcc/tree-core.h for the new uses of TREE_READONLY/readonly_flag.
>
> We basically all use OMP_CLAUSE_SUBCODE_CHECK macros for OpenMP clause expressions exclusively,
> so I don't see a reason to diverge from that style (even when context is clear).

ACK.

> I have greatly expanded the test scan patterns to include parallel/kernels/serial/data/enter data,
> as well as non-readonly copyin clause together with readonly.

Thanks.

> Also added simple 'declare' tests, but there is not anything to scan in the 'tree-original' dump though.

Yeah, the current OpenACC 'declare' implementation is "special".

>>> --- 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'.
>
> Fixed in this v3 patch.

Thanks.

> Again, tested on x86_64-linux + nvptx offloading. Okay for mainline?

Yes, thanks.


Grüße
 Thomas


> gcc/c/ChangeLog:
>
> 	* c-parser.cc (c_parser_oacc_data_clause): Add parsing support for
> 	'readonly' modifier, set OMP_CLAUSE_MAP_READONLY if readonly modifier
> 	found, update comments.
> 	(c_parser_oacc_cache): Add parsing support for 'readonly' modifier,
> 	set OMP_CLAUSE__CACHE__READONLY if readonly modifier found, update
> 	comments.
>
> gcc/cp/ChangeLog:
>
> 	* parser.cc (cp_parser_oacc_data_clause): Add parsing support for
> 	'readonly' modifier, set OMP_CLAUSE_MAP_READONLY if readonly modifier
> 	found, update comments.
> 	(cp_parser_oacc_cache): Add parsing support for 'readonly' modifier,
> 	set OMP_CLAUSE__CACHE__READONLY if readonly modifier found, update
> 	comments.
>
> gcc/fortran/ChangeLog:
>
> 	* dump-parse-tree.cc (show_omp_namelist): Print "readonly," for
> 	OMP_LIST_MAP and OMP_LIST_CACHE if n->u.map.readonly is set.
> 	Adjust 'n->u.map_op' to 'n->u.map.op'.
> 	* gfortran.h (typedef struct gfc_omp_namelist): Adjust map_op as
> 	'ENUM_BITFIELD (gfc_omp_map_op) op:8', add 'bool readonly' field,
> 	change to named struct field 'map'.
>
> 	* openmp.cc (gfc_match_omp_map_clause): Adjust 'n->u.map_op' to
> 	'n->u.map.op'.
> 	(gfc_match_omp_clause_reduction): Likewise.
>
> 	(gfc_match_omp_clauses): Add readonly modifier parsing for OpenACC
> 	copyin clause, set 'n->u.map.op' and 'n->u.map.readonly' for parsed
> 	clause. Adjust 'n->u.map_op' to 'n->u.map.op'.
> 	(gfc_match_oacc_declare): Adjust 'n->u.map_op' to 'n->u.map.op'.
> 	(gfc_match_oacc_cache): Add readonly modifier parsing for OpenACC
> 	cache directive.
> 	(resolve_omp_clauses): Adjust 'n->u.map_op' to 'n->u.map.op'.
> 	* trans-decl.cc (add_clause): Adjust 'n->u.map_op' to 'n->u.map.op'.
> 	(finish_oacc_declare): Likewise.
> 	* trans-openmp.cc (gfc_trans_omp_clauses): Set OMP_CLAUSE_MAP_READONLY,
> 	OMP_CLAUSE__CACHE__READONLY to 1 when readonly is set. Adjust
> 	'n->u.map_op' to 'n->u.map.op'.
> 	(gfc_add_clause_implicitly): Adjust 'n->u.map_op' to 'n->u.map.op'.
>
> gcc/ChangeLog:
> 	* tree.h (OMP_CLAUSE_MAP_READONLY): New macro.
> 	(OMP_CLAUSE__CACHE__READONLY): New macro.
> 	* tree-core.h (struct GTY(()) tree_base): Adjust comments for new
> 	uses of readonly_flag bit in OMP_CLAUSE_MAP_READONLY and
> 	OMP_CLAUSE__CACHE__READONLY.
> 	* tree-pretty-print.cc (dump_omp_clause): Add support for printing
> 	OMP_CLAUSE_MAP_READONLY and OMP_CLAUSE__CACHE__READONLY.
>
> gcc/testsuite/ChangeLog:
>
> 	* c-c++-common/goacc/readonly-1.c: New test.
> 	* gfortran.dg/goacc/readonly-1.f90: New test.
>
>
>
>
>
> diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
> index 53e99aa29d9..00f8bf4376e 100644
> --- a/gcc/c/c-parser.cc
> +++ b/gcc/c/c-parser.cc
> @@ -15627,7 +15627,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,
> @@ -15680,11 +15684,37 @@ 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, false);
>  
> -  for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
> -    OMP_CLAUSE_SET_MAP_KIND (c, kind);
> +  tree nl = list;
> +  bool readonly = false;
> +  location_t open_loc = c_parser_peek_token (parser)->location;
> +  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;
> +	    }
> +	}
> +      nl = c_parser_omp_variable_list (parser, open_loc, OMP_CLAUSE_MAP, list,
> +				       false);
> +      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;
>  }
> @@ -19821,15 +19851,39 @@ 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;
> +  location_t open_loc = c_parser_peek_token (parser)->location;
> +  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;
> +	}
> +      clauses = c_parser_omp_variable_list (parser, open_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);
> diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
> index e32acfc30a2..4fe27fb07b2 100644
> --- a/gcc/cp/parser.cc
> +++ b/gcc/cp/parser.cc
> @@ -38544,7 +38544,11 @@ cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list,
>     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
>  cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
> @@ -38597,11 +38601,34 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
>      default:
>        gcc_unreachable ();
>      }
> -  tree nl, c;
> -  nl = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list, false);
>  
> -  for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
> -    OMP_CLAUSE_SET_MAP_KIND (c, kind);
> +  tree nl = list;
> +  bool readonly = false;
> +  if (cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
> +    {
> +      /* Turn on readonly modifier parsing for copyin clause.  */
> +      if (c_kind == PRAGMA_OACC_CLAUSE_COPYIN)
> +	{
> +	  cp_token *token = cp_lexer_peek_token (parser->lexer);
> +	  if (token->type == CPP_NAME
> +	      && !strcmp (IDENTIFIER_POINTER (token->u.value), "readonly")
> +	      && cp_lexer_peek_nth_token (parser->lexer, 2)->type == CPP_COLON)
> +	    {
> +	      cp_lexer_consume_token (parser->lexer);
> +	      cp_lexer_consume_token (parser->lexer);
> +	      readonly = true;
> +	    }
> +	}
> +      nl = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE_MAP, list, NULL,
> +					   false);
> +    }
> +
> +  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;
>  }
> @@ -47178,6 +47205,9 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
>  
>  /* OpenACC 2.0:
>     # pragma acc cache (variable-list) new-line
> +
> +   OpenACC 2.7:
> +   # pragma acc cache (readonly: variable-list) new-line
>  */
>  
>  static tree
> @@ -47187,9 +47217,28 @@ cp_parser_oacc_cache (cp_parser *parser, cp_token *pragma_tok)
>       clauses.  */
>    auto_suppress_location_wrappers sentinel;
>  
> -  tree stmt, clauses;
> +  tree stmt, clauses = NULL_TREE;
> +  bool readonly = false;
> +
> +  if (cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
> +    {
> +      cp_token *token = cp_lexer_peek_token (parser->lexer);
> +      if (token->type == CPP_NAME
> +	  && !strcmp (IDENTIFIER_POINTER (token->u.value), "readonly")
> +	  && cp_lexer_peek_nth_token (parser->lexer, 2)->type == CPP_COLON)
> +	{
> +	  cp_lexer_consume_token (parser->lexer);
> +	  cp_lexer_consume_token (parser->lexer);
> +	  readonly = true;
> +	}
> +      clauses = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE__CACHE_,
> +						NULL, NULL);
> +    }
> +
> +  if (readonly)
> +    for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
> +      OMP_CLAUSE__CACHE__READONLY (c) = 1;
>  
> -  clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE__CACHE_, NULL_TREE);
>    clauses = finish_omp_clauses (clauses, C_ORT_ACC);
>  
>    cp_parser_require_pragma_eol (parser, cp_lexer_peek_token (parser->lexer));
> diff --git a/gcc/fortran/dump-parse-tree.cc b/gcc/fortran/dump-parse-tree.cc
> index 7b154eb3ca7..db84b06289b 100644
> --- a/gcc/fortran/dump-parse-tree.cc
> +++ b/gcc/fortran/dump-parse-tree.cc
> @@ -1400,6 +1400,9 @@ show_omp_namelist (int list_type, gfc_omp_namelist *n)
>  	    fputs (") ALLOCATE(", dumpfile);
>  	  continue;
>  	}
> +      if ((list_type == OMP_LIST_MAP || list_type == OMP_LIST_CACHE)
> +	  && n->u.map.readonly)
> +	fputs ("readonly,", dumpfile);
>        if (list_type == OMP_LIST_REDUCTION)
>  	switch (n->u.reduction_op)
>  	  {
> @@ -1467,7 +1470,7 @@ show_omp_namelist (int list_type, gfc_omp_namelist *n)
>  	  default: break;
>  	  }
>        else if (list_type == OMP_LIST_MAP)
> -	switch (n->u.map_op)
> +	switch (n->u.map.op)
>  	  {
>  	  case OMP_MAP_ALLOC: fputs ("alloc:", dumpfile); break;
>  	  case OMP_MAP_TO: fputs ("to:", dumpfile); break;
> diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
> index ebba2336e12..32b792f85fb 100644
> --- a/gcc/fortran/gfortran.h
> +++ b/gcc/fortran/gfortran.h
> @@ -1363,7 +1363,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) op:8;
> +	  bool readonly;
> +        } map;
>        gfc_expr *align;
>        struct
>  	{
> diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
> index 38de60238c0..5c44e666eb9 100644
> --- a/gcc/fortran/openmp.cc
> +++ b/gcc/fortran/openmp.cc
> @@ -1210,7 +1210,7 @@ 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;
>        return true;
>      }
>  
> @@ -1524,7 +1524,7 @@ gfc_match_omp_clause_reduction (char pc, gfc_omp_clauses *c, bool openacc,
>  	    gfc_omp_namelist *p = gfc_get_omp_namelist (), **tl;
>  	    p->sym = n->sym;
>  	    p->where = p->where;
> -	    p->u.map_op = OMP_MAP_ALWAYS_TOFROM;
> +	    p->u.map.op = OMP_MAP_ALWAYS_TOFROM;
>  
>  	    tl = &c->lists[OMP_LIST_MAP];
>  	    while (*tl)
> @@ -2181,11 +2181,25 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
>  	    {
>  	      if (openacc)
>  		{
> -		  if (gfc_match ("copyin ( ") == MATCH_YES
> -		      && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
> -						   OMP_MAP_TO, true,
> -						   allow_derived))
> -		    continue;
> +		  if (gfc_match ("copyin ( ") == MATCH_YES)
> +		    {
> +		      bool readonly = gfc_match ("readonly : ") == MATCH_YES;
> +		      head = NULL;
> +		      if (gfc_match_omp_variable_list ("",
> +						       &c->lists[OMP_LIST_MAP],
> +						       true, NULL, &head, true,
> +						       allow_derived)
> +			  == MATCH_YES)
> +			{
> +			  gfc_omp_namelist *n;
> +			  for (n = *head; n; n = n->next)
> +			    {
> +			      n->u.map.op = OMP_MAP_TO;
> +			      n->u.map.readonly = readonly;
> +			    }
> +			  continue;
> +			}
> +		    }
>  		}
>  	      else if (gfc_match_omp_variable_list ("copyin (",
>  						    &c->lists[OMP_LIST_COPYIN],
> @@ -3134,7 +3148,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
>  		{
>  		  gfc_omp_namelist *n;
>  		  for (n = *head; n; n = n->next)
> -		    n->u.map_op = map_op;
> +		    n->u.map.op = map_op;
>  		  continue;
>  		}
>  	      gfc_current_locus = old_loc;
> @@ -4002,7 +4016,7 @@ gfc_match_oacc_declare (void)
>        if (gfc_current_ns->proc_name
>  	  && gfc_current_ns->proc_name->attr.flavor == FL_MODULE)
>  	{
> -	  if (n->u.map_op != OMP_MAP_ALLOC && n->u.map_op != OMP_MAP_TO)
> +	  if (n->u.map.op != OMP_MAP_ALLOC && n->u.map.op != OMP_MAP_TO)
>  	    {
>  	      gfc_error ("Invalid clause in module with !$ACC DECLARE at %L",
>  			 &where);
> @@ -4036,7 +4050,7 @@ gfc_match_oacc_declare (void)
>  	  return MATCH_ERROR;
>  	}
>  
> -      switch (n->u.map_op)
> +      switch (n->u.map.op)
>  	{
>  	  case OMP_MAP_FORCE_ALLOC:
>  	  case OMP_MAP_ALLOC:
> @@ -4151,21 +4165,36 @@ gfc_match_oacc_wait (void)
>  match
>  gfc_match_oacc_cache (void)
>  {
> +  bool readonly = false;
>    gfc_omp_clauses *c = gfc_get_omp_clauses ();
>    /* The OpenACC cache directive explicitly only allows "array elements or
>       subarrays", which we're currently not checking here.  Either check this
>       after the call of gfc_match_omp_variable_list, or add something like a
>       only_sections variant next to its allow_sections parameter.  */
> -  match m = gfc_match_omp_variable_list (" (",
> -					 &c->lists[OMP_LIST_CACHE], true,
> -					 NULL, NULL, true);
> +  match m = gfc_match (" ( ");
>    if (m != MATCH_YES)
>      {
>        gfc_free_omp_clauses(c);
>        return m;
>      }
>  
> -  if (gfc_current_state() != COMP_DO 
> +  if (gfc_match ("readonly : ") == MATCH_YES)
> +    readonly = true;
> +
> +  gfc_omp_namelist **head = NULL;
> +  m = gfc_match_omp_variable_list ("", &c->lists[OMP_LIST_CACHE], true,
> +				   NULL, &head, true);
> +  if (m != MATCH_YES)
> +    {
> +      gfc_free_omp_clauses(c);
> +      return m;
> +    }
> +
> +  if (readonly)
> +    for (gfc_omp_namelist *n = *head; n; n = n->next)
> +      n->u.map.readonly = true;
> +
> +  if (gfc_current_state() != COMP_DO
>        && gfc_current_state() != COMP_DO_CONCURRENT)
>      {
>        gfc_error ("ACC CACHE directive must be inside of loop %C");
> @@ -8436,8 +8465,8 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
>  		  }
>  		if (openacc
>  		    && list == OMP_LIST_MAP
> -		    && (n->u.map_op == OMP_MAP_ATTACH
> -			|| n->u.map_op == OMP_MAP_DETACH))
> +		    && (n->u.map.op == OMP_MAP_ATTACH
> +			|| n->u.map.op == OMP_MAP_DETACH))
>  		  {
>  		    symbol_attribute attr;
>  		    if (n->expr)
> @@ -8447,7 +8476,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
>  		    if (!attr.pointer && !attr.allocatable)
>  		      gfc_error ("%qs clause argument must be ALLOCATABLE or "
>  				 "a POINTER at %L",
> -				 (n->u.map_op == OMP_MAP_ATTACH) ? "attach"
> +				 (n->u.map.op == OMP_MAP_ATTACH) ? "attach"
>  				 : "detach", &n->where);
>  		  }
>  		if (lastref
> @@ -8518,7 +8547,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
>  		else if (openacc)
>  		  {
>  		    if (list == OMP_LIST_MAP
> -			&& n->u.map_op == OMP_MAP_FORCE_DEVICEPTR)
> +			&& n->u.map.op == OMP_MAP_FORCE_DEVICEPTR)
>  		      resolve_oacc_deviceptr_clause (n->sym, n->where, name);
>  		    else
>  		      resolve_oacc_data_clauses (n->sym, n->where, name);
> @@ -8540,7 +8569,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
>  		    {
>  		    case EXEC_OMP_TARGET:
>  		    case EXEC_OMP_TARGET_DATA:
> -		      switch (n->u.map_op)
> +		      switch (n->u.map.op)
>  			{
>  			case OMP_MAP_TO:
>  			case OMP_MAP_ALWAYS_TO:
> @@ -8567,7 +8596,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
>  			}
>  		      break;
>  		    case EXEC_OMP_TARGET_ENTER_DATA:
> -		      switch (n->u.map_op)
> +		      switch (n->u.map.op)
>  			{
>  			case OMP_MAP_TO:
>  			case OMP_MAP_ALWAYS_TO:
> @@ -8577,16 +8606,16 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
>  			case OMP_MAP_PRESENT_ALLOC:
>  			  break;
>  			case OMP_MAP_TOFROM:
> -			  n->u.map_op = OMP_MAP_TO;
> +			  n->u.map.op = OMP_MAP_TO;
>  			  break;
>  			case OMP_MAP_ALWAYS_TOFROM:
> -			  n->u.map_op = OMP_MAP_ALWAYS_TO;
> +			  n->u.map.op = OMP_MAP_ALWAYS_TO;
>  			  break;
>  			case OMP_MAP_PRESENT_TOFROM:
> -			  n->u.map_op = OMP_MAP_PRESENT_TO;
> +			  n->u.map.op = OMP_MAP_PRESENT_TO;
>  			  break;
>  			case OMP_MAP_ALWAYS_PRESENT_TOFROM:
> -			  n->u.map_op = OMP_MAP_ALWAYS_PRESENT_TO;
> +			  n->u.map.op = OMP_MAP_ALWAYS_PRESENT_TO;
>  			  break;
>  			default:
>  			  gfc_error ("TARGET ENTER DATA with map-type other "
> @@ -8596,7 +8625,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
>  			}
>  		      break;
>  		    case EXEC_OMP_TARGET_EXIT_DATA:
> -		      switch (n->u.map_op)
> +		      switch (n->u.map.op)
>  			{
>  			case OMP_MAP_FROM:
>  			case OMP_MAP_ALWAYS_FROM:
> @@ -8606,16 +8635,16 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
>  			case OMP_MAP_DELETE:
>  			  break;
>  			case OMP_MAP_TOFROM:
> -			  n->u.map_op = OMP_MAP_FROM;
> +			  n->u.map.op = OMP_MAP_FROM;
>  			  break;
>  			case OMP_MAP_ALWAYS_TOFROM:
> -			  n->u.map_op = OMP_MAP_ALWAYS_FROM;
> +			  n->u.map.op = OMP_MAP_ALWAYS_FROM;
>  			  break;
>  			case OMP_MAP_PRESENT_TOFROM:
> -			  n->u.map_op = OMP_MAP_PRESENT_FROM;
> +			  n->u.map.op = OMP_MAP_PRESENT_FROM;
>  			  break;
>  			case OMP_MAP_ALWAYS_PRESENT_TOFROM:
> -			  n->u.map_op = OMP_MAP_ALWAYS_PRESENT_FROM;
> +			  n->u.map.op = OMP_MAP_ALWAYS_PRESENT_FROM;
>  			  break;
>  			default:
>  			  gfc_error ("TARGET EXIT DATA with map-type other "
> diff --git a/gcc/fortran/trans-decl.cc b/gcc/fortran/trans-decl.cc
> index 6d463036966..b7dea11461f 100644
> --- a/gcc/fortran/trans-decl.cc
> +++ b/gcc/fortran/trans-decl.cc
> @@ -6744,7 +6744,7 @@ add_clause (gfc_symbol *sym, gfc_omp_map_op map_op)
>  
>    n = gfc_get_omp_namelist ();
>    n->sym = sym;
> -  n->u.map_op = map_op;
> +  n->u.map.op = map_op;
>  
>    if (!module_oacc_clauses)
>      module_oacc_clauses = gfc_get_omp_clauses ();
> @@ -6846,10 +6846,10 @@ finish_oacc_declare (gfc_namespace *ns, gfc_symbol *sym, bool block)
>  
>    for (n = omp_clauses->lists[OMP_LIST_MAP]; n; n = n->next)
>      {
> -      switch (n->u.map_op)
> +      switch (n->u.map.op)
>  	{
>  	  case OMP_MAP_DEVICE_RESIDENT:
> -	    n->u.map_op = OMP_MAP_FORCE_ALLOC;
> +	    n->u.map.op = OMP_MAP_FORCE_ALLOC;
>  	    break;
>  
>  	  default:
> diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
> index a2bf15665b3..fa1bfd41380 100644
> --- a/gcc/fortran/trans-openmp.cc
> +++ b/gcc/fortran/trans-openmp.cc
> @@ -3139,7 +3139,10 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
>  		      || (n->expr && gfc_expr_attr (n->expr).pointer)))
>  		always_modifier = true;
>  
> -	      switch (n->u.map_op)
> +	      if (n->u.map.readonly)
> +		OMP_CLAUSE_MAP_READONLY (node) = 1;
> +
> +	      switch (n->u.map.op)
>  		{
>  		case OMP_MAP_ALLOC:
>  		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC);
> @@ -3266,8 +3269,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
>  		      && n->sym->attr.omp_declare_target
>  		      && (always_modifier || n->sym->attr.pointer)
>  		      && op != EXEC_OMP_TARGET_EXIT_DATA
> -		      && n->u.map_op != OMP_MAP_DELETE
> -		      && n->u.map_op != OMP_MAP_RELEASE)
> +		      && n->u.map.op != OMP_MAP_DELETE
> +		      && n->u.map.op != OMP_MAP_RELEASE)
>  		    {
>  		      gcc_assert (n->sym->ts.u.cl->backend_decl);
>  		      node5 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
> @@ -3333,7 +3336,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
>  			{
>  			  enum gomp_map_kind gmk = GOMP_MAP_POINTER;
>  			  if (op == EXEC_OMP_TARGET_EXIT_DATA
> -			      && n->u.map_op == OMP_MAP_DELETE)
> +			      && n->u.map.op == OMP_MAP_DELETE)
>  			    gmk = GOMP_MAP_DELETE;
>  			  else if (op == EXEC_OMP_TARGET_EXIT_DATA)
>  			    gmk = GOMP_MAP_RELEASE;
> @@ -3356,7 +3359,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
>  			{
>  			  enum gomp_map_kind gmk;
>  			  if (op == EXEC_OMP_TARGET_EXIT_DATA
> -			      && n->u.map_op == OMP_MAP_DELETE)
> +			      && n->u.map.op == OMP_MAP_DELETE)
>  			    gmk = GOMP_MAP_DELETE;
>  			  else if (op == EXEC_OMP_TARGET_EXIT_DATA)
>  			    gmk = GOMP_MAP_RELEASE;
> @@ -3388,18 +3391,18 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
>  		      node2 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
>  		      OMP_CLAUSE_DECL (node2) = decl;
>  		      OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
> -		      if (n->u.map_op == OMP_MAP_DELETE)
> +		      if (n->u.map.op == OMP_MAP_DELETE)
>  			map_kind = GOMP_MAP_DELETE;
>  		      else if (op == EXEC_OMP_TARGET_EXIT_DATA
> -			       || n->u.map_op == OMP_MAP_RELEASE)
> +			       || n->u.map.op == OMP_MAP_RELEASE)
>  			map_kind = GOMP_MAP_RELEASE;
>  		      else
>  			map_kind = GOMP_MAP_TO_PSET;
>  		      OMP_CLAUSE_SET_MAP_KIND (node2, map_kind);
>  
>  		      if (op != EXEC_OMP_TARGET_EXIT_DATA
> -			  && n->u.map_op != OMP_MAP_DELETE
> -			  && n->u.map_op != OMP_MAP_RELEASE)
> +			  && n->u.map.op != OMP_MAP_DELETE
> +			  && n->u.map.op != OMP_MAP_RELEASE)
>  			{
>  			  node3 = build_omp_clause (input_location,
>  						    OMP_CLAUSE_MAP);
> @@ -3417,7 +3420,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
>  			      = gfc_conv_descriptor_data_get (decl);
>  			  OMP_CLAUSE_SIZE (node3) = size_int (0);
>  
> -			  if (n->u.map_op == OMP_MAP_ATTACH)
> +			  if (n->u.map.op == OMP_MAP_ATTACH)
>  			    {
>  			      /* Standalone attach clauses used with arrays with
>  				 descriptors must copy the descriptor to the
> @@ -3433,7 +3436,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
>  			      node3 = NULL;
>  			      goto finalize_map_clause;
>  			    }
> -			  else if (n->u.map_op == OMP_MAP_DETACH)
> +			  else if (n->u.map.op == OMP_MAP_DETACH)
>  			    {
>  			      OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_DETACH);
>  			      /* Similarly to above, we don't want to unmap PTR
> @@ -3626,8 +3629,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
>  			 to perform a single attach/detach operation, of the
>  			 pointer itself, not of the pointed-to object.  */
>  		      if (openacc
> -			  && (n->u.map_op == OMP_MAP_ATTACH
> -			      || n->u.map_op == OMP_MAP_DETACH))
> +			  && (n->u.map.op == OMP_MAP_ATTACH
> +			      || n->u.map.op == OMP_MAP_DETACH))
>  			{
>  			  OMP_CLAUSE_DECL (node)
>  			    = build_fold_addr_expr (OMP_CLAUSE_DECL (node));
> @@ -3656,7 +3659,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
>  					       se.string_length),
>  					   TYPE_SIZE_UNIT (tmp));
>  			  gomp_map_kind kind;
> -			  if (n->u.map_op == OMP_MAP_DELETE)
> +			  if (n->u.map.op == OMP_MAP_DELETE)
>  			    kind = GOMP_MAP_DELETE;
>  			  else if (op == EXEC_OMP_TARGET_EXIT_DATA)
>  			    kind = GOMP_MAP_RELEASE;
> @@ -3713,8 +3716,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
>  			     to perform a single attach/detach operation, of the
>  			     pointer itself, not of the pointed-to object.  */
>  			  if (openacc
> -			      && (n->u.map_op == OMP_MAP_ATTACH
> -				  || n->u.map_op == OMP_MAP_DETACH))
> +			      && (n->u.map.op == OMP_MAP_ATTACH
> +				  || n->u.map.op == OMP_MAP_DETACH))
>  			    {
>  			      OMP_CLAUSE_DECL (node)
>  				= build_fold_addr_expr (inner);
> @@ -3806,8 +3809,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
>  		    {
>  		      /* Bare attach and detach clauses don't want any
>  			 additional nodes.  */
> -		      if ((n->u.map_op == OMP_MAP_ATTACH
> -			   || n->u.map_op == OMP_MAP_DETACH)
> +		      if ((n->u.map.op == OMP_MAP_ATTACH
> +			   || n->u.map.op == OMP_MAP_DETACH)
>  			  && (POINTER_TYPE_P (TREE_TYPE (inner))
>  			      || GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (inner))))
>  			{
> @@ -3840,8 +3843,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
>  			    map_kind = ((GOMP_MAP_ALWAYS_P (map_kind)
>  					 || gfc_expr_attr (n->expr).pointer)
>  					? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO);
> -			  else if (n->u.map_op == OMP_MAP_RELEASE
> -				   || n->u.map_op == OMP_MAP_DELETE)
> +			  else if (n->u.map.op == OMP_MAP_RELEASE
> +				   || n->u.map.op == OMP_MAP_DELETE)
>  			    ;
>  			  else if (op == EXEC_OMP_TARGET_EXIT_DATA
>  				   || op == EXEC_OACC_EXIT_DATA)
> @@ -4088,6 +4091,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
>  		}
>  	      if (n->u.present_modifier)
>  		OMP_CLAUSE_MOTION_PRESENT (node) = 1;
> +	      if (list == OMP_LIST_CACHE && n->u.map.readonly)
> +		OMP_CLAUSE__CACHE__READONLY (node) = 1;
>  	      omp_clauses = gfc_trans_add_clause (node, omp_clauses);
>  	    }
>  	  break;
> @@ -6561,7 +6566,7 @@ gfc_add_clause_implicitly (gfc_omp_clauses *clauses_out,
>  	  n2->where = n->where;
>  	  n2->sym = n->sym;
>  	  if (is_target)
> -	    n2->u.map_op = OMP_MAP_TOFROM;
> +	    n2->u.map.op = OMP_MAP_TOFROM;
>  	  if (tail)
>  	    {
>  	      tail->next = n2;
> diff --git a/gcc/testsuite/c-c++-common/goacc/readonly-1.c b/gcc/testsuite/c-c++-common/goacc/readonly-1.c
> new file mode 100644
> index 00000000000..34fc92c24d5
> --- /dev/null
> +++ b/gcc/testsuite/c-c++-common/goacc/readonly-1.c
> @@ -0,0 +1,59 @@
> +/* { dg-additional-options "-fdump-tree-original" } */
> +
> +struct S
> +{
> +  int *ptr;
> +  float f;
> +};
> +
> +int a[32], b[32];
> +#pragma acc declare copyin(readonly: a) copyin(b)
> +
> +int main (void)
> +{
> +  int x[32], y[32];
> +  struct S s = {x, 0};
> +
> +  #pragma acc parallel copyin(readonly: x[:32], s.ptr[:16]) copyin(y[:32])
> +  {
> +    #pragma acc cache (readonly: x[:32])
> +    #pragma acc cache (y[:32])
> +  }
> +
> +  #pragma acc kernels copyin(readonly: x[:32], s.ptr[:16]) copyin(y[:32])
> +  {
> +    #pragma acc cache (readonly: x[:32])
> +    #pragma acc cache (y[:32])
> +  }
> +
> +  #pragma acc serial copyin(readonly: x[:32], s.ptr[:16]) copyin(y[:32])
> +  {
> +    #pragma acc cache (readonly: x[:32])
> +    #pragma acc cache (y[:32])
> +  }
> +
> +  #pragma acc data copyin(readonly: x[:32], s.ptr[:16]) copyin(y[:32])
> +  {
> +    #pragma acc cache (readonly: x[:32])
> +    #pragma acc cache (y[:32])
> +  }
> +
> +  #pragma acc enter data copyin(readonly: x[:32], s.ptr[:16]) copyin(y[:32])
> +
> +  return 0;
> +}
> +
> +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */
> +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */
> +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */
> +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */
> +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */
> +
> +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */
> +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */
> +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */
> +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */
> +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */
> +
> +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(readonly:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\);$" 4 "original" } } */
> +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\);$" 4 "original" } } */
> diff --git a/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90 b/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90
> new file mode 100644
> index 00000000000..696ebd08321
> --- /dev/null
> +++ b/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90
> @@ -0,0 +1,89 @@
> +! { dg-additional-options "-fdump-tree-original" }
> +
> +subroutine foo (a, n)
> +  integer :: n, a(:)
> +  integer :: i, b(n), c(n)
> +  !$acc parallel copyin(readonly: a(:), b(:n)) copyin(c(:))
> +  do i = 1,32
> +     !$acc cache (readonly: a(:), b(:n))
> +     !$acc cache (c(:))
> +  enddo
> +  !$acc end parallel
> +
> +  !$acc kernels copyin(readonly: a(:), b(:n)) copyin(c(:))
> +  do i = 1,32
> +     !$acc cache (readonly: a(:), b(:n))
> +     !$acc cache (c(:))
> +  enddo
> +  !$acc end kernels
> +
> +  !$acc serial copyin(readonly: a(:), b(:n)) copyin(c(:))
> +  do i = 1,32
> +     !$acc cache (readonly: a(:), b(:n))
> +     !$acc cache (c(:))
> +  enddo
> +  !$acc end serial
> +
> +  !$acc data copyin(readonly: a(:), b(:n)) copyin(c(:))
> +  do i = 1,32
> +     !$acc cache (readonly: a(:), b(:n))
> +     !$acc cache (c(:))
> +  enddo
> +  !$acc end data
> +
> +  !$acc enter data copyin(readonly: a(:), b(:n)) copyin(c(:))
> +
> +end subroutine foo
> +
> +program main
> +  integer :: g(32), h(32)
> +  integer :: i, n = 32, a(32)
> +  integer :: b(32), c(32)
> +
> +  !$acc declare copyin(readonly: g), copyin(h)
> +
> +  !$acc parallel copyin(readonly: a(:32), b(:n)) copyin(c(:))
> +  do i = 1,32
> +     !$acc cache (readonly: a(:), b(:n))
> +     !$acc cache (c(:))
> +  enddo
> +  !$acc end parallel
> +
> +  !$acc kernels copyin(readonly: a(:), b(:n)) copyin(c(:))
> +  do i = 1,32
> +     !$acc cache (readonly: a(:), b(:n))
> +     !$acc cache (c(:))
> +  enddo
> +  !$acc end kernels
> +
> +  !$acc serial copyin(readonly: a(:), b(:n)) copyin(c(:))
> +  do i = 1,32
> +     !$acc cache (readonly: a(:), b(:n))
> +     !$acc cache (c(:))
> +  enddo
> +  !$acc end serial
> +
> +  !$acc data copyin(readonly: a(:), b(:n)) copyin(c(:))
> +  do i = 1,32
> +     !$acc cache (readonly: a(:), b(:n))
> +     !$acc cache (c(:))
> +  enddo
> +  !$acc end data
> +
> +  !$acc enter data copyin(readonly: a(:), b(:n)) copyin(c(:))
> +
> +end program main
> +
> +! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } }
> +! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } }
> +! { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(readonly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } }
> +! { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(readonly,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } }
> +! { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(readonly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } }
> +! { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(readonly,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } }
> +! { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(readonly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } }
> +! { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(readonly,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } }
> +! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(readonly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } }
> +! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(readonly,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } }
> +
> +! { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(readonly:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\) \\(readonly:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\);" 8 "original" } }
> +! { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\);" 8 "original" } }
> diff --git a/gcc/tree-core.h b/gcc/tree-core.h
> index 8a89462bd7e..d529712306d 100644
> --- a/gcc/tree-core.h
> +++ b/gcc/tree-core.h
> @@ -1344,6 +1344,12 @@ struct GTY(()) tree_base {
>         TYPE_READONLY in
>             all types
>  
> +       OMP_CLAUSE_MAP_READONLY in
> +           OMP_CLAUSE_MAP
> +
> +       OMP_CLAUSE__CACHE__READONLY in
> +           OMP_CLAUSE__CACHE_
> +
>     constant_flag:
>  
>         TREE_CONSTANT in
> diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
> index 654f5247e3a..926f7e006a7 100644
> --- a/gcc/tree-pretty-print.cc
> +++ b/gcc/tree-pretty-print.cc
> @@ -913,6 +913,8 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
>  
>      case OMP_CLAUSE_MAP:
>        pp_string (pp, "map(");
> +      if (OMP_CLAUSE_MAP_READONLY (clause))
> +	pp_string (pp, "readonly,");
>        switch (OMP_CLAUSE_MAP_KIND (clause))
>  	{
>  	case GOMP_MAP_ALLOC:
> @@ -1095,6 +1097,8 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
>  
>      case OMP_CLAUSE__CACHE_:
>        pp_string (pp, "(");
> +      if (OMP_CLAUSE__CACHE__READONLY (clause))
> +	pp_string (pp, "readonly:");
>        dump_generic_node (pp, OMP_CLAUSE_DECL (clause),
>  			 spc, flags, false);
>        goto print_clause_size;
> diff --git a/gcc/tree.h b/gcc/tree.h
> index e1fc6c2221d..b67a37d6522 100644
> --- a/gcc/tree.h
> +++ b/gcc/tree.h
> @@ -1841,6 +1841,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_))
> +
>  /* True on an OMP_CLAUSE_USE_DEVICE_PTR with an OpenACC 'if_present'
>     clause.  */
>  #define OMP_CLAUSE_USE_DEVICE_PTR_IF_PRESENT(NODE) \

OpenACC 2.7: front-end support for readonly modifier: Add basic OpenACC 'declare' testing (was: [PATCH, OpenACC 2.7, v2] readonly modifier support in front-ends)

[Thomas Schwinge]

[-- Attachment #1: Type: text/plain, Size: 23603 bytes --]

Hi!

On 2024-03-13T10:12:17+0100, I wrote:
> On 2024-03-07T17:02:02+0900, Chung-Lin Tang <cltang@pllab.cs.nthu.edu.tw> wrote:
>> Also added simple 'declare' tests, but there is not anything to scan in the 'tree-original' dump though.
>
> Yeah, the current OpenACC 'declare' implementation is "special".

Actually -- commit 38958ac987dc3e6162e2ddaba3c7e7f41381e079
"OpenACC 2.7: front-end support for readonly modifier: Add basic OpenACC 'declare' testing",
see attached.


But I realized another thing: don't we have to handle the 'readonly'
modifier also in Fortran module files, that is, next to the OpenACC
'declare' 'copyin' handling in 'gcc/fortran/module.cc':
'AB_OACC_DECLARE_COPYIN' etc.?  Chung-Lin, please check, via test cases.
'gfortran.dg/goacc/routine-module*', for example, should provide some
guidance of how to achieve actual module file use, and then do the same
'scan-tree-dump' as in the current 'readonly' modifier test cases.
I suppose the code changes would look similar to
commit a61f6afbee370785cf091fe46e2e022748528307
"OpenACC 'nohost' clause", for example.  By means of only emitting a tag
in the module file if the 'readonly' modifier is specified, we should
maintain compatibility with the current 'MOD_VERSION'.


Grüße
 Thomas


>> diff --git a/gcc/fortran/dump-parse-tree.cc b/gcc/fortran/dump-parse-tree.cc
>> index 7b154eb3ca7..db84b06289b 100644
>> --- a/gcc/fortran/dump-parse-tree.cc
>> +++ b/gcc/fortran/dump-parse-tree.cc
>> @@ -1400,6 +1400,9 @@ show_omp_namelist (int list_type, gfc_omp_namelist *n)
>>  	    fputs (") ALLOCATE(", dumpfile);
>>  	  continue;
>>  	}
>> +      if ((list_type == OMP_LIST_MAP || list_type == OMP_LIST_CACHE)
>> +	  && n->u.map.readonly)
>> +	fputs ("readonly,", dumpfile);
>>        if (list_type == OMP_LIST_REDUCTION)
>>  	switch (n->u.reduction_op)
>>  	  {
>> @@ -1467,7 +1470,7 @@ show_omp_namelist (int list_type, gfc_omp_namelist *n)
>>  	  default: break;
>>  	  }
>>        else if (list_type == OMP_LIST_MAP)
>> -	switch (n->u.map_op)
>> +	switch (n->u.map.op)
>>  	  {
>>  	  case OMP_MAP_ALLOC: fputs ("alloc:", dumpfile); break;
>>  	  case OMP_MAP_TO: fputs ("to:", dumpfile); break;
>> diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
>> index ebba2336e12..32b792f85fb 100644
>> --- a/gcc/fortran/gfortran.h
>> +++ b/gcc/fortran/gfortran.h
>> @@ -1363,7 +1363,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) op:8;
>> +	  bool readonly;
>> +        } map;
>>        gfc_expr *align;
>>        struct
>>  	{
>> diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
>> index 38de60238c0..5c44e666eb9 100644
>> --- a/gcc/fortran/openmp.cc
>> +++ b/gcc/fortran/openmp.cc
>> @@ -1210,7 +1210,7 @@ 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;
>>        return true;
>>      }
>>  
>> @@ -1524,7 +1524,7 @@ gfc_match_omp_clause_reduction (char pc, gfc_omp_clauses *c, bool openacc,
>>  	    gfc_omp_namelist *p = gfc_get_omp_namelist (), **tl;
>>  	    p->sym = n->sym;
>>  	    p->where = p->where;
>> -	    p->u.map_op = OMP_MAP_ALWAYS_TOFROM;
>> +	    p->u.map.op = OMP_MAP_ALWAYS_TOFROM;
>>  
>>  	    tl = &c->lists[OMP_LIST_MAP];
>>  	    while (*tl)
>> @@ -2181,11 +2181,25 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
>>  	    {
>>  	      if (openacc)
>>  		{
>> -		  if (gfc_match ("copyin ( ") == MATCH_YES
>> -		      && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
>> -						   OMP_MAP_TO, true,
>> -						   allow_derived))
>> -		    continue;
>> +		  if (gfc_match ("copyin ( ") == MATCH_YES)
>> +		    {
>> +		      bool readonly = gfc_match ("readonly : ") == MATCH_YES;
>> +		      head = NULL;
>> +		      if (gfc_match_omp_variable_list ("",
>> +						       &c->lists[OMP_LIST_MAP],
>> +						       true, NULL, &head, true,
>> +						       allow_derived)
>> +			  == MATCH_YES)
>> +			{
>> +			  gfc_omp_namelist *n;
>> +			  for (n = *head; n; n = n->next)
>> +			    {
>> +			      n->u.map.op = OMP_MAP_TO;
>> +			      n->u.map.readonly = readonly;
>> +			    }
>> +			  continue;
>> +			}
>> +		    }
>>  		}
>>  	      else if (gfc_match_omp_variable_list ("copyin (",
>>  						    &c->lists[OMP_LIST_COPYIN],
>> @@ -3134,7 +3148,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
>>  		{
>>  		  gfc_omp_namelist *n;
>>  		  for (n = *head; n; n = n->next)
>> -		    n->u.map_op = map_op;
>> +		    n->u.map.op = map_op;
>>  		  continue;
>>  		}
>>  	      gfc_current_locus = old_loc;
>> @@ -4002,7 +4016,7 @@ gfc_match_oacc_declare (void)
>>        if (gfc_current_ns->proc_name
>>  	  && gfc_current_ns->proc_name->attr.flavor == FL_MODULE)
>>  	{
>> -	  if (n->u.map_op != OMP_MAP_ALLOC && n->u.map_op != OMP_MAP_TO)
>> +	  if (n->u.map.op != OMP_MAP_ALLOC && n->u.map.op != OMP_MAP_TO)
>>  	    {
>>  	      gfc_error ("Invalid clause in module with !$ACC DECLARE at %L",
>>  			 &where);
>> @@ -4036,7 +4050,7 @@ gfc_match_oacc_declare (void)
>>  	  return MATCH_ERROR;
>>  	}
>>  
>> -      switch (n->u.map_op)
>> +      switch (n->u.map.op)
>>  	{
>>  	  case OMP_MAP_FORCE_ALLOC:
>>  	  case OMP_MAP_ALLOC:
>> @@ -4151,21 +4165,36 @@ gfc_match_oacc_wait (void)
>>  match
>>  gfc_match_oacc_cache (void)
>>  {
>> +  bool readonly = false;
>>    gfc_omp_clauses *c = gfc_get_omp_clauses ();
>>    /* The OpenACC cache directive explicitly only allows "array elements or
>>       subarrays", which we're currently not checking here.  Either check this
>>       after the call of gfc_match_omp_variable_list, or add something like a
>>       only_sections variant next to its allow_sections parameter.  */
>> -  match m = gfc_match_omp_variable_list (" (",
>> -					 &c->lists[OMP_LIST_CACHE], true,
>> -					 NULL, NULL, true);
>> +  match m = gfc_match (" ( ");
>>    if (m != MATCH_YES)
>>      {
>>        gfc_free_omp_clauses(c);
>>        return m;
>>      }
>>  
>> -  if (gfc_current_state() != COMP_DO 
>> +  if (gfc_match ("readonly : ") == MATCH_YES)
>> +    readonly = true;
>> +
>> +  gfc_omp_namelist **head = NULL;
>> +  m = gfc_match_omp_variable_list ("", &c->lists[OMP_LIST_CACHE], true,
>> +				   NULL, &head, true);
>> +  if (m != MATCH_YES)
>> +    {
>> +      gfc_free_omp_clauses(c);
>> +      return m;
>> +    }
>> +
>> +  if (readonly)
>> +    for (gfc_omp_namelist *n = *head; n; n = n->next)
>> +      n->u.map.readonly = true;
>> +
>> +  if (gfc_current_state() != COMP_DO
>>        && gfc_current_state() != COMP_DO_CONCURRENT)
>>      {
>>        gfc_error ("ACC CACHE directive must be inside of loop %C");
>> @@ -8436,8 +8465,8 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
>>  		  }
>>  		if (openacc
>>  		    && list == OMP_LIST_MAP
>> -		    && (n->u.map_op == OMP_MAP_ATTACH
>> -			|| n->u.map_op == OMP_MAP_DETACH))
>> +		    && (n->u.map.op == OMP_MAP_ATTACH
>> +			|| n->u.map.op == OMP_MAP_DETACH))
>>  		  {
>>  		    symbol_attribute attr;
>>  		    if (n->expr)
>> @@ -8447,7 +8476,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
>>  		    if (!attr.pointer && !attr.allocatable)
>>  		      gfc_error ("%qs clause argument must be ALLOCATABLE or "
>>  				 "a POINTER at %L",
>> -				 (n->u.map_op == OMP_MAP_ATTACH) ? "attach"
>> +				 (n->u.map.op == OMP_MAP_ATTACH) ? "attach"
>>  				 : "detach", &n->where);
>>  		  }
>>  		if (lastref
>> @@ -8518,7 +8547,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
>>  		else if (openacc)
>>  		  {
>>  		    if (list == OMP_LIST_MAP
>> -			&& n->u.map_op == OMP_MAP_FORCE_DEVICEPTR)
>> +			&& n->u.map.op == OMP_MAP_FORCE_DEVICEPTR)
>>  		      resolve_oacc_deviceptr_clause (n->sym, n->where, name);
>>  		    else
>>  		      resolve_oacc_data_clauses (n->sym, n->where, name);
>> @@ -8540,7 +8569,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
>>  		    {
>>  		    case EXEC_OMP_TARGET:
>>  		    case EXEC_OMP_TARGET_DATA:
>> -		      switch (n->u.map_op)
>> +		      switch (n->u.map.op)
>>  			{
>>  			case OMP_MAP_TO:
>>  			case OMP_MAP_ALWAYS_TO:
>> @@ -8567,7 +8596,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
>>  			}
>>  		      break;
>>  		    case EXEC_OMP_TARGET_ENTER_DATA:
>> -		      switch (n->u.map_op)
>> +		      switch (n->u.map.op)
>>  			{
>>  			case OMP_MAP_TO:
>>  			case OMP_MAP_ALWAYS_TO:
>> @@ -8577,16 +8606,16 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
>>  			case OMP_MAP_PRESENT_ALLOC:
>>  			  break;
>>  			case OMP_MAP_TOFROM:
>> -			  n->u.map_op = OMP_MAP_TO;
>> +			  n->u.map.op = OMP_MAP_TO;
>>  			  break;
>>  			case OMP_MAP_ALWAYS_TOFROM:
>> -			  n->u.map_op = OMP_MAP_ALWAYS_TO;
>> +			  n->u.map.op = OMP_MAP_ALWAYS_TO;
>>  			  break;
>>  			case OMP_MAP_PRESENT_TOFROM:
>> -			  n->u.map_op = OMP_MAP_PRESENT_TO;
>> +			  n->u.map.op = OMP_MAP_PRESENT_TO;
>>  			  break;
>>  			case OMP_MAP_ALWAYS_PRESENT_TOFROM:
>> -			  n->u.map_op = OMP_MAP_ALWAYS_PRESENT_TO;
>> +			  n->u.map.op = OMP_MAP_ALWAYS_PRESENT_TO;
>>  			  break;
>>  			default:
>>  			  gfc_error ("TARGET ENTER DATA with map-type other "
>> @@ -8596,7 +8625,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
>>  			}
>>  		      break;
>>  		    case EXEC_OMP_TARGET_EXIT_DATA:
>> -		      switch (n->u.map_op)
>> +		      switch (n->u.map.op)
>>  			{
>>  			case OMP_MAP_FROM:
>>  			case OMP_MAP_ALWAYS_FROM:
>> @@ -8606,16 +8635,16 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
>>  			case OMP_MAP_DELETE:
>>  			  break;
>>  			case OMP_MAP_TOFROM:
>> -			  n->u.map_op = OMP_MAP_FROM;
>> +			  n->u.map.op = OMP_MAP_FROM;
>>  			  break;
>>  			case OMP_MAP_ALWAYS_TOFROM:
>> -			  n->u.map_op = OMP_MAP_ALWAYS_FROM;
>> +			  n->u.map.op = OMP_MAP_ALWAYS_FROM;
>>  			  break;
>>  			case OMP_MAP_PRESENT_TOFROM:
>> -			  n->u.map_op = OMP_MAP_PRESENT_FROM;
>> +			  n->u.map.op = OMP_MAP_PRESENT_FROM;
>>  			  break;
>>  			case OMP_MAP_ALWAYS_PRESENT_TOFROM:
>> -			  n->u.map_op = OMP_MAP_ALWAYS_PRESENT_FROM;
>> +			  n->u.map.op = OMP_MAP_ALWAYS_PRESENT_FROM;
>>  			  break;
>>  			default:
>>  			  gfc_error ("TARGET EXIT DATA with map-type other "
>> diff --git a/gcc/fortran/trans-decl.cc b/gcc/fortran/trans-decl.cc
>> index 6d463036966..b7dea11461f 100644
>> --- a/gcc/fortran/trans-decl.cc
>> +++ b/gcc/fortran/trans-decl.cc
>> @@ -6744,7 +6744,7 @@ add_clause (gfc_symbol *sym, gfc_omp_map_op map_op)
>>  
>>    n = gfc_get_omp_namelist ();
>>    n->sym = sym;
>> -  n->u.map_op = map_op;
>> +  n->u.map.op = map_op;
>>  
>>    if (!module_oacc_clauses)
>>      module_oacc_clauses = gfc_get_omp_clauses ();
>> @@ -6846,10 +6846,10 @@ finish_oacc_declare (gfc_namespace *ns, gfc_symbol *sym, bool block)
>>  
>>    for (n = omp_clauses->lists[OMP_LIST_MAP]; n; n = n->next)
>>      {
>> -      switch (n->u.map_op)
>> +      switch (n->u.map.op)
>>  	{
>>  	  case OMP_MAP_DEVICE_RESIDENT:
>> -	    n->u.map_op = OMP_MAP_FORCE_ALLOC;
>> +	    n->u.map.op = OMP_MAP_FORCE_ALLOC;
>>  	    break;
>>  
>>  	  default:
>> diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
>> index a2bf15665b3..fa1bfd41380 100644
>> --- a/gcc/fortran/trans-openmp.cc
>> +++ b/gcc/fortran/trans-openmp.cc
>> @@ -3139,7 +3139,10 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
>>  		      || (n->expr && gfc_expr_attr (n->expr).pointer)))
>>  		always_modifier = true;
>>  
>> -	      switch (n->u.map_op)
>> +	      if (n->u.map.readonly)
>> +		OMP_CLAUSE_MAP_READONLY (node) = 1;
>> +
>> +	      switch (n->u.map.op)
>>  		{
>>  		case OMP_MAP_ALLOC:
>>  		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC);
>> @@ -3266,8 +3269,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
>>  		      && n->sym->attr.omp_declare_target
>>  		      && (always_modifier || n->sym->attr.pointer)
>>  		      && op != EXEC_OMP_TARGET_EXIT_DATA
>> -		      && n->u.map_op != OMP_MAP_DELETE
>> -		      && n->u.map_op != OMP_MAP_RELEASE)
>> +		      && n->u.map.op != OMP_MAP_DELETE
>> +		      && n->u.map.op != OMP_MAP_RELEASE)
>>  		    {
>>  		      gcc_assert (n->sym->ts.u.cl->backend_decl);
>>  		      node5 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
>> @@ -3333,7 +3336,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
>>  			{
>>  			  enum gomp_map_kind gmk = GOMP_MAP_POINTER;
>>  			  if (op == EXEC_OMP_TARGET_EXIT_DATA
>> -			      && n->u.map_op == OMP_MAP_DELETE)
>> +			      && n->u.map.op == OMP_MAP_DELETE)
>>  			    gmk = GOMP_MAP_DELETE;
>>  			  else if (op == EXEC_OMP_TARGET_EXIT_DATA)
>>  			    gmk = GOMP_MAP_RELEASE;
>> @@ -3356,7 +3359,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
>>  			{
>>  			  enum gomp_map_kind gmk;
>>  			  if (op == EXEC_OMP_TARGET_EXIT_DATA
>> -			      && n->u.map_op == OMP_MAP_DELETE)
>> +			      && n->u.map.op == OMP_MAP_DELETE)
>>  			    gmk = GOMP_MAP_DELETE;
>>  			  else if (op == EXEC_OMP_TARGET_EXIT_DATA)
>>  			    gmk = GOMP_MAP_RELEASE;
>> @@ -3388,18 +3391,18 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
>>  		      node2 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
>>  		      OMP_CLAUSE_DECL (node2) = decl;
>>  		      OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
>> -		      if (n->u.map_op == OMP_MAP_DELETE)
>> +		      if (n->u.map.op == OMP_MAP_DELETE)
>>  			map_kind = GOMP_MAP_DELETE;
>>  		      else if (op == EXEC_OMP_TARGET_EXIT_DATA
>> -			       || n->u.map_op == OMP_MAP_RELEASE)
>> +			       || n->u.map.op == OMP_MAP_RELEASE)
>>  			map_kind = GOMP_MAP_RELEASE;
>>  		      else
>>  			map_kind = GOMP_MAP_TO_PSET;
>>  		      OMP_CLAUSE_SET_MAP_KIND (node2, map_kind);
>>  
>>  		      if (op != EXEC_OMP_TARGET_EXIT_DATA
>> -			  && n->u.map_op != OMP_MAP_DELETE
>> -			  && n->u.map_op != OMP_MAP_RELEASE)
>> +			  && n->u.map.op != OMP_MAP_DELETE
>> +			  && n->u.map.op != OMP_MAP_RELEASE)
>>  			{
>>  			  node3 = build_omp_clause (input_location,
>>  						    OMP_CLAUSE_MAP);
>> @@ -3417,7 +3420,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
>>  			      = gfc_conv_descriptor_data_get (decl);
>>  			  OMP_CLAUSE_SIZE (node3) = size_int (0);
>>  
>> -			  if (n->u.map_op == OMP_MAP_ATTACH)
>> +			  if (n->u.map.op == OMP_MAP_ATTACH)
>>  			    {
>>  			      /* Standalone attach clauses used with arrays with
>>  				 descriptors must copy the descriptor to the
>> @@ -3433,7 +3436,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
>>  			      node3 = NULL;
>>  			      goto finalize_map_clause;
>>  			    }
>> -			  else if (n->u.map_op == OMP_MAP_DETACH)
>> +			  else if (n->u.map.op == OMP_MAP_DETACH)
>>  			    {
>>  			      OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_DETACH);
>>  			      /* Similarly to above, we don't want to unmap PTR
>> @@ -3626,8 +3629,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
>>  			 to perform a single attach/detach operation, of the
>>  			 pointer itself, not of the pointed-to object.  */
>>  		      if (openacc
>> -			  && (n->u.map_op == OMP_MAP_ATTACH
>> -			      || n->u.map_op == OMP_MAP_DETACH))
>> +			  && (n->u.map.op == OMP_MAP_ATTACH
>> +			      || n->u.map.op == OMP_MAP_DETACH))
>>  			{
>>  			  OMP_CLAUSE_DECL (node)
>>  			    = build_fold_addr_expr (OMP_CLAUSE_DECL (node));
>> @@ -3656,7 +3659,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
>>  					       se.string_length),
>>  					   TYPE_SIZE_UNIT (tmp));
>>  			  gomp_map_kind kind;
>> -			  if (n->u.map_op == OMP_MAP_DELETE)
>> +			  if (n->u.map.op == OMP_MAP_DELETE)
>>  			    kind = GOMP_MAP_DELETE;
>>  			  else if (op == EXEC_OMP_TARGET_EXIT_DATA)
>>  			    kind = GOMP_MAP_RELEASE;
>> @@ -3713,8 +3716,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
>>  			     to perform a single attach/detach operation, of the
>>  			     pointer itself, not of the pointed-to object.  */
>>  			  if (openacc
>> -			      && (n->u.map_op == OMP_MAP_ATTACH
>> -				  || n->u.map_op == OMP_MAP_DETACH))
>> +			      && (n->u.map.op == OMP_MAP_ATTACH
>> +				  || n->u.map.op == OMP_MAP_DETACH))
>>  			    {
>>  			      OMP_CLAUSE_DECL (node)
>>  				= build_fold_addr_expr (inner);
>> @@ -3806,8 +3809,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
>>  		    {
>>  		      /* Bare attach and detach clauses don't want any
>>  			 additional nodes.  */
>> -		      if ((n->u.map_op == OMP_MAP_ATTACH
>> -			   || n->u.map_op == OMP_MAP_DETACH)
>> +		      if ((n->u.map.op == OMP_MAP_ATTACH
>> +			   || n->u.map.op == OMP_MAP_DETACH)
>>  			  && (POINTER_TYPE_P (TREE_TYPE (inner))
>>  			      || GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (inner))))
>>  			{
>> @@ -3840,8 +3843,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
>>  			    map_kind = ((GOMP_MAP_ALWAYS_P (map_kind)
>>  					 || gfc_expr_attr (n->expr).pointer)
>>  					? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO);
>> -			  else if (n->u.map_op == OMP_MAP_RELEASE
>> -				   || n->u.map_op == OMP_MAP_DELETE)
>> +			  else if (n->u.map.op == OMP_MAP_RELEASE
>> +				   || n->u.map.op == OMP_MAP_DELETE)
>>  			    ;
>>  			  else if (op == EXEC_OMP_TARGET_EXIT_DATA
>>  				   || op == EXEC_OACC_EXIT_DATA)
>> @@ -4088,6 +4091,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
>>  		}
>>  	      if (n->u.present_modifier)
>>  		OMP_CLAUSE_MOTION_PRESENT (node) = 1;
>> +	      if (list == OMP_LIST_CACHE && n->u.map.readonly)
>> +		OMP_CLAUSE__CACHE__READONLY (node) = 1;
>>  	      omp_clauses = gfc_trans_add_clause (node, omp_clauses);
>>  	    }
>>  	  break;
>> @@ -6561,7 +6566,7 @@ gfc_add_clause_implicitly (gfc_omp_clauses *clauses_out,
>>  	  n2->where = n->where;
>>  	  n2->sym = n->sym;
>>  	  if (is_target)
>> -	    n2->u.map_op = OMP_MAP_TOFROM;
>> +	    n2->u.map.op = OMP_MAP_TOFROM;
>>  	  if (tail)
>>  	    {
>>  	      tail->next = n2;

>> diff --git a/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90 b/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90
>> new file mode 100644
>> index 00000000000..696ebd08321
>> --- /dev/null
>> +++ b/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90
>> @@ -0,0 +1,89 @@
>> +! { dg-additional-options "-fdump-tree-original" }
>> +
>> +subroutine foo (a, n)
>> +  integer :: n, a(:)
>> +  integer :: i, b(n), c(n)
>> +  !$acc parallel copyin(readonly: a(:), b(:n)) copyin(c(:))
>> +  do i = 1,32
>> +     !$acc cache (readonly: a(:), b(:n))
>> +     !$acc cache (c(:))
>> +  enddo
>> +  !$acc end parallel
>> +
>> +  !$acc kernels copyin(readonly: a(:), b(:n)) copyin(c(:))
>> +  do i = 1,32
>> +     !$acc cache (readonly: a(:), b(:n))
>> +     !$acc cache (c(:))
>> +  enddo
>> +  !$acc end kernels
>> +
>> +  !$acc serial copyin(readonly: a(:), b(:n)) copyin(c(:))
>> +  do i = 1,32
>> +     !$acc cache (readonly: a(:), b(:n))
>> +     !$acc cache (c(:))
>> +  enddo
>> +  !$acc end serial
>> +
>> +  !$acc data copyin(readonly: a(:), b(:n)) copyin(c(:))
>> +  do i = 1,32
>> +     !$acc cache (readonly: a(:), b(:n))
>> +     !$acc cache (c(:))
>> +  enddo
>> +  !$acc end data
>> +
>> +  !$acc enter data copyin(readonly: a(:), b(:n)) copyin(c(:))
>> +
>> +end subroutine foo
>> +
>> +program main
>> +  integer :: g(32), h(32)
>> +  integer :: i, n = 32, a(32)
>> +  integer :: b(32), c(32)
>> +
>> +  !$acc declare copyin(readonly: g), copyin(h)
>> +
>> +  !$acc parallel copyin(readonly: a(:32), b(:n)) copyin(c(:))
>> +  do i = 1,32
>> +     !$acc cache (readonly: a(:), b(:n))
>> +     !$acc cache (c(:))
>> +  enddo
>> +  !$acc end parallel
>> +
>> +  !$acc kernels copyin(readonly: a(:), b(:n)) copyin(c(:))
>> +  do i = 1,32
>> +     !$acc cache (readonly: a(:), b(:n))
>> +     !$acc cache (c(:))
>> +  enddo
>> +  !$acc end kernels
>> +
>> +  !$acc serial copyin(readonly: a(:), b(:n)) copyin(c(:))
>> +  do i = 1,32
>> +     !$acc cache (readonly: a(:), b(:n))
>> +     !$acc cache (c(:))
>> +  enddo
>> +  !$acc end serial
>> +
>> +  !$acc data copyin(readonly: a(:), b(:n)) copyin(c(:))
>> +  do i = 1,32
>> +     !$acc cache (readonly: a(:), b(:n))
>> +     !$acc cache (c(:))
>> +  enddo
>> +  !$acc end data
>> +
>> +  !$acc enter data copyin(readonly: a(:), b(:n)) copyin(c(:))
>> +
>> +end program main
>> +
>> +! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } }
>> +! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } }
>> +! { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(readonly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } }
>> +! { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(readonly,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } }
>> +! { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(readonly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } }
>> +! { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(readonly,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } }
>> +! { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(readonly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } }
>> +! { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(readonly,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } }
>> +! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(readonly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } }
>> +! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(readonly,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } }
>> +
>> +! { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(readonly:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\) \\(readonly:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\);" 8 "original" } }
>> +! { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\);" 8 "original" } }



[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-OpenACC-2.7-front-end-support-for-readonly-modifier-.patch --]
[-- Type: text/x-diff, Size: 3887 bytes --]

From 38958ac987dc3e6162e2ddaba3c7e7f41381e079 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tschwinge@baylibre.com>
Date: Thu, 14 Mar 2024 15:01:01 +0100
Subject: [PATCH] OpenACC 2.7: front-end support for readonly modifier: Add
 basic OpenACC 'declare' testing

... to complement commit ddf852dac2abaca317c10b8323f338123b0585c8
"OpenACC 2.7: front-end support for readonly modifier".

	gcc/testsuite/
	* c-c++-common/goacc/readonly-1.c: Add basic OpenACC 'declare'
	testing.
	* gfortran.dg/goacc/readonly-1.f90: Likewise.
---
 gcc/testsuite/c-c++-common/goacc/readonly-1.c  | 5 +++++
 gcc/testsuite/gfortran.dg/goacc/readonly-1.f90 | 6 ++++++
 2 files changed, 11 insertions(+)

diff --git a/gcc/testsuite/c-c++-common/goacc/readonly-1.c b/gcc/testsuite/c-c++-common/goacc/readonly-1.c
index 34fc92c24d5..300464c92e3 100644
--- a/gcc/testsuite/c-c++-common/goacc/readonly-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/readonly-1.c
@@ -8,12 +8,15 @@ struct S
 
 int a[32], b[32];
 #pragma acc declare copyin(readonly: a) copyin(b)
+/* Not visible in 'original' dump; handled via 'offload_vars'.  */
 
 int main (void)
 {
   int x[32], y[32];
   struct S s = {x, 0};
 
+  #pragma acc declare copyin(readonly: x/*[:32]*/, s/*.ptr[:16]*/) copyin(y/*[:32]*/)
+
   #pragma acc parallel copyin(readonly: x[:32], s.ptr[:16]) copyin(y[:32])
   {
     #pragma acc cache (readonly: x[:32])
@@ -43,6 +46,8 @@ int main (void)
   return 0;
 }
 
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc declare map\\(to:y\\) map\\(readonly,to:s\\) map\\(readonly,to:x\\)" 1 "original" } } */
+
 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */
 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */
 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(to:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "original" { target { c } } } } */
diff --git a/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90 b/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90
index 696ebd08321..fc1e2719e67 100644
--- a/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90
@@ -3,6 +3,9 @@
 subroutine foo (a, n)
   integer :: n, a(:)
   integer :: i, b(n), c(n)
+  !!$acc declare copyin(readonly: a(:), b(:n)) copyin(c(:))
+  !$acc declare copyin(readonly: b) copyin(c)
+
   !$acc parallel copyin(readonly: a(:), b(:n)) copyin(c(:))
   do i = 1,32
      !$acc cache (readonly: a(:), b(:n))
@@ -74,6 +77,9 @@ program main
 
 end program main
 
+! The front end turns OpenACC 'declare' into OpenACC 'data'.
+!   { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(readonly,to:\\*b\\) map\\(alloc:b.+ map\\(to:\\*c\\) map\\(alloc:c.+" 1 "original" } }
+!   { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(readonly,to:g\\) map\\(to:h\\)" 1 "original" } }
 ! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } }
 ! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(to:c.+ map\\(alloc:c.+" 1 "original" } }
 ! { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(readonly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(to:\\*.+ map\\(alloc:c.+" 1 "original" } }
-- 
2.34.1

Re: OpenACC 2.7: front-end support for readonly modifier: Add basic OpenACC 'declare' testing

[Tobias Burnus]

[-- Attachment #1: Type: text/plain, Size: 1351 bytes --]

Hi all, hi Thomas & Chung-Lin,

Thomas Schwinge wrote:
> But I realized another thing: don't we have to handle the 'readonly'
> modifier also in Fortran module files, that is, next to the OpenACC
> 'declare' 'copyin' handling in 'gcc/fortran/module.cc':
> 'AB_OACC_DECLARE_COPYIN' etc.?

I bet so; it is not as bad as with the others as it is "only" an
optimization hint, but it makes sense to make it available.

Note that when you place the 'module' in the same file as the module
users ('use'), the compiler might know things because they are in the
same translation unit / file not because it is in the module ...

>   Chung-Lin, please check, via test cases.
> 'gfortran.dg/goacc/routine-module*', for example, should provide some
> guidance of how to achieve actual module file use, and then do the same
> 'scan-tree-dump' as in the current 'readonly' modifier test cases.
...
> By means of only emitting a tag
> in the module file if the 'readonly' modifier is specified, we should
> maintain compatibility with the current 'MOD_VERSION'.

That was the idea: If only new information gets added (if used), older
compilers still work. This has huge limitations and does not work as
well as imagined but here it should work: Older .mod will work with new
compilers, even though the reverse might not be true.

Tobias

Re: OpenACC 2.7: front-end support for readonly modifier: Add basic OpenACC 'declare' testing

[Tobias Burnus]

[-- Attachment #1: Type: text/plain, Size: 1327 bytes --]

Hi all, hi Thomas & Chung-Lin,

Thomas Schwinge wrote:
> But I realized another thing: don't we have to handle the 'readonly'
> modifier also in Fortran module files, that is, next to the OpenACC
> 'declare' 'copyin' handling in 'gcc/fortran/module.cc':
> 'AB_OACC_DECLARE_COPYIN' etc.?

I bet so; it is not as bad as with the others as it is "only" an 
optimization hint, but it makes sense to make it available.

Note that when you place the 'module' in the same file as the module 
users ('use'), the compiler might know things because they are in the 
same translation unit / file not because it is in the module ...

>   Chung-Lin, please check, via test cases.
> 'gfortran.dg/goacc/routine-module*', for example, should provide some
> guidance of how to achieve actual module file use, and then do the same
> 'scan-tree-dump' as in the current 'readonly' modifier test cases.
...
> By means of only emitting a tag
> in the module file if the 'readonly' modifier is specified, we should
> maintain compatibility with the current 'MOD_VERSION'.

That was the idea: If only new information gets added (if used), older 
compilers still work. This has huge limitations and does not work as 
well as imagined but here it should work: Older .mod will work with new 
compilers, even though the reverse might not be true.

Tobias

Re: [PATCH, OpenACC 2.7] Connect readonly modifier to points-to analysis

[Thomas Schwinge]

Hi!

Richard, as the original author of 'SSA_NAME_POINTS_TO_READONLY_MEMORY':
2018 commit 6214d5c7e7470bdd5ecbeae668c2522551bfebbc (Subversion r263958)
"Move const_parm trick to generic code"; 'gcc/tree.h':

    /* Nonzero if this SSA_NAME is known to point to memory that may not
       be written to.  This is set for default defs of function parameters
       that have a corresponding r or R specification in the functions
       fn spec attribute.  This is used by alias analysis.  */
    #define SSA_NAME_POINTS_TO_READONLY_MEMORY(NODE) \
        SSA_NAME_CHECK (NODE)->base.deprecated_flag

..., may I ask you to please help review the following patch
(full-quoted)?

For context: this patch here ("second patch") depends on a first patch:
<inbox.sourceware.org/d0e6013f-ca38-b98d-dc01-b30adbd5901a@siemens.com>
"[PATCH, OpenACC 2.7] readonly modifier support in front-ends".  That one
is still under review/rework; so you're not able to apply this second
patch here.

In a nutshell: a 'readonly' modifier has been added to the OpenACC
'copyin' clause (copy host to device memory, don't copy back at end of
region):

| If the optional 'readonly' modifier appears, then the implementation may assume that the data
| referenced by _var-list_ is never written to within the applicable region.

That is, for example (untested):

    #pragma acc routine
    void escape(int *);

    int x[32] = [...];
    #pragma acc parallel copyin(readonly: x)
    {
      int a1 = x[3];
      escape(x);
      int a2 = x[3]; // Per 'readonly', don't need to reload 'x[3]' here.
      //x[22] = 0; // Invalid -- but no diagnostic mandated.
    }

What Chung-Lin's first patch does is mark the OMP clause for 'x' (not the
'x' decl itself!) as 'readonly', via a new 'OMP_CLAUSE_MAP_READONLY'
flag.

The actual optimization then is done in this second patch.  Chung-Lin
found that he could use 'SSA_NAME_POINTS_TO_READONLY_MEMORY' for that.
I don't have much experience with most of the following generic code, so
would appreciate a helping hand, whether that conceptually makes sense as
well as from the implementation point of view:

On 2023-07-25T23:52:06+0800, Chung-Lin Tang via Gcc-patches <gcc-patches@gcc.gnu.org> wrote:
> On 2023/7/11 2:33 AM, Chung-Lin Tang via Gcc-patches wrote:
>> As we discussed earlier, the work for actually linking this to middle-end
>> points-to analysis is a somewhat non-trivial issue. This first patch allows
>> the language feature to be used in OpenACC directives first (with no effect for now).
>> The middle-end changes are probably going to be a later patch.
>
> This second patch tries to link the readonly modifier to points-to analysis.
>
> There already exists SSA_NAME_POINTS_TO_READONLY_MEMORY and it's support in the
> alias oracle routines in tree-ssa-alias.cc, so basically what this patch does is
> try to make the variables holding the array section base pointers to have this
> flag set.
>
> There is an another OMP_CLAUSE_MAP_POINTS_TO_READONLY set by front-ends on the
> associated pointer clauses if OMP_CLAUSE_MAP_READONLY is set.
> Also a DECL_POINTS_TO_READONLY flag is set for VAR_DECLs when creating the tmp
> vars carrying these receiver references on the offloaded side. These
> eventually get translated to SSA_NAME_POINTS_TO_READONLY_MEMORY.


> This still doesn't always work as expected in terms of optimization:
> struct pointer fields and Fortran arrays (kind of like C structs) which have
> several accesses to create the pointer access on the receive/offloaded side,
> and SRA appears to not work on these sequences, so gets in the way of much
> redundancy elimination.

I understand correctly that this is left as future work?  Please add the test
cases you have, XFAILed in some reasonable way.


> Currently have one testcase where we can demonstrate 'readonly' can avoid
> a clobber by function call.

:-)


> --- a/gcc/c/c-typeck.cc
> +++ b/gcc/c/c-typeck.cc
> @@ -14258,6 +14258,8 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
>       OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
>        else
>       OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
> +      if (OMP_CLAUSE_MAP_READONLY (c))
> +     OMP_CLAUSE_MAP_POINTS_TO_READONLY (c2) = 1;
>        OMP_CLAUSE_MAP_IMPLICIT (c2) = OMP_CLAUSE_MAP_IMPLICIT (c);
>        if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
>         && !c_mark_addressable (t))

> --- a/gcc/cp/semantics.cc
> +++ b/gcc/cp/semantics.cc
> @@ -5872,6 +5872,8 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
>           }
>         else
>           OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
> +       if (OMP_CLAUSE_MAP_READONLY (c))
> +         OMP_CLAUSE_MAP_POINTS_TO_READONLY (c2) = 1;
>         OMP_CLAUSE_MAP_IMPLICIT (c2) = OMP_CLAUSE_MAP_IMPLICIT (c);
>         if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
>             && !cxx_mark_addressable (t))

> --- a/gcc/fortran/trans-openmp.cc
> +++ b/gcc/fortran/trans-openmp.cc
> @@ -2524,6 +2524,8 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_exec_op op,
>        node3 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
>        OMP_CLAUSE_SET_MAP_KIND (node3, ptr_kind);
>        OMP_CLAUSE_DECL (node3) = gfc_conv_descriptor_data_get (decl);
> +      if (n->u.readonly)
> +     OMP_CLAUSE_MAP_POINTS_TO_READONLY (node3) = 1;
>        /* This purposely does not include GOMP_MAP_ALWAYS_POINTER.  The extra
>        cast prevents gimplify.cc from recognising it as being part of the
>        struct - and adding an 'alloc: for the 'desc.data' pointer, which
> @@ -2559,6 +2561,8 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_exec_op op,
>                               OMP_CLAUSE_MAP);
>        OMP_CLAUSE_SET_MAP_KIND (node3, ptr_kind);
>        OMP_CLAUSE_DECL (node3) = decl;
> +      if (n->u.readonly)
> +     OMP_CLAUSE_MAP_POINTS_TO_READONLY (node3) = 1;
>      }

Could combine these two into one, after
'if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))' reconverges here, like
where 'OMP_CLAUSE_SIZE (node3)' is set:

>    ptr2 = fold_convert (ptrdiff_type_node, ptr2);
>    OMP_CLAUSE_SIZE (node3) = fold_build2 (MINUS_EXPR, ptrdiff_type_node,

Is 'n->u.readonly == OMP_CLAUSE_MAP_READONLY (node)'?  If yes, would the
latter be clearer to use as the 'if' expression (like in C, C++ front
ends)?

I see further additional 'OMP_CLAUSE_MAP' clauses synthesized, for
example in 'gcc/cp/semantics.cc:handle_omp_array_sections', or
'gcc/fortran/trans-openmp.cc:gfc_trans_omp_array_section', also
'gcc/gimplify.cc'.  I assume these are not relevant to have
'OMP_CLAUSE_MAP_READONLY' -> 'OMP_CLAUSE_MAP_POINTS_TO_READONLY'
propagated?  Actually, per your changes (see below), there is one
'OMP_CLAUSE_MAP_POINTS_TO_READONLY' propagation in
'gcc/gimplify.cc:build_omp_struct_comp_nodes'.

Is the current situation re flag setting/propagation what was empirically
necessary to make the test case work, or is it a systematic review?  (The
former is fine; I'd just like to know.)

> --- a/gcc/gimple-expr.cc
> +++ b/gcc/gimple-expr.cc
> @@ -376,6 +376,8 @@ copy_var_decl (tree var, tree name, tree type)
>    DECL_CONTEXT (copy) = DECL_CONTEXT (var);
>    TREE_USED (copy) = 1;
>    DECL_SEEN_IN_BIND_EXPR_P (copy) = 1;
> +  if (VAR_P (var))
> +    DECL_POINTS_TO_READONLY (copy) = DECL_POINTS_TO_READONLY (var);
>    DECL_ATTRIBUTES (copy) = DECL_ATTRIBUTES (var);
>    if (DECL_USER_ALIGN (var))
>      {

> --- a/gcc/gimplify.cc
> +++ b/gcc/gimplify.cc
> @@ -221,6 +221,7 @@ struct gimplify_omp_ctx
>    splay_tree variables;
>    hash_set<tree> *privatized_types;
>    tree clauses;
> +  hash_set<tree_operand_hash> *pt_readonly_ptrs;
>    /* Iteration variables in an OMP_FOR.  */
>    vec<tree> loop_iter_var;
>    location_t location;
> @@ -628,6 +629,15 @@ internal_get_tmp_var (tree val, gimple_seq *pre_p, gimple_seq *post_p,
>    gimplify_expr (&val, pre_p, post_p, is_gimple_reg_rhs_or_call,
>                fb_rvalue);
>
> +  bool pt_readonly = false;
> +  if (gimplify_omp_ctxp && gimplify_omp_ctxp->pt_readonly_ptrs)
> +    {
> +      tree ptr = val;
> +      if (TREE_CODE (ptr) == POINTER_PLUS_EXPR)
> +     ptr = TREE_OPERAND (ptr, 0);
> +      pt_readonly = gimplify_omp_ctxp->pt_readonly_ptrs->contains (ptr);
> +    }

'POINTER_PLUS_EXPR' is the only special thing we may run into, here?
(Generally, I prefer 'if', 'else if, [...], 'else gcc_unreachable ()'.)

> +
>    if (allow_ssa
>        && gimplify_ctxp->into_ssa
>        && is_gimple_reg_type (TREE_TYPE (val)))
> @@ -639,9 +649,18 @@ internal_get_tmp_var (tree val, gimple_seq *pre_p, gimple_seq *post_p,
>         if (name)
>           SET_SSA_NAME_VAR_OR_IDENTIFIER (t, create_tmp_var_name (name));
>       }
> +      if (pt_readonly)
> +     SSA_NAME_POINTS_TO_READONLY_MEMORY (t) = 1;
>      }
>    else
> -    t = lookup_tmp_var (val, is_formal, not_gimple_reg);
> +    {
> +      t = lookup_tmp_var (val, is_formal, not_gimple_reg);
> +      if (pt_readonly)
> +     {
> +       DECL_POINTS_TO_READONLY (t) = 1;
> +       gimplify_omp_ctxp->pt_readonly_ptrs->add (t);
> +     }
> +    }
>
>    mod = build2 (INIT_EXPR, TREE_TYPE (t), t, unshare_expr (val));
>
> @@ -8906,6 +8925,8 @@ build_omp_struct_comp_nodes (enum tree_code code, tree grp_start, tree grp_end,
>    OMP_CLAUSE_SET_MAP_KIND (c2, mkind);
>    OMP_CLAUSE_DECL (c2) = unshare_expr (OMP_CLAUSE_DECL (grp_end));
>    OMP_CLAUSE_CHAIN (c2) = NULL_TREE;
> +  if (OMP_CLAUSE_MAP_POINTS_TO_READONLY (grp_end))
> +    OMP_CLAUSE_MAP_POINTS_TO_READONLY (c2) = 1;
>    tree grp_mid = NULL_TREE;
>    if (OMP_CLAUSE_CHAIN (grp_start) != grp_end)
>      grp_mid = OMP_CLAUSE_CHAIN (grp_start);

For my understanding, is this empirically necessary, or a systematic
review?

> @@ -11741,6 +11762,16 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
>
>             gimplify_omp_ctxp = outer_ctx;
>           }
> +       else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
> +                && (code == OACC_PARALLEL
> +                    || code == OACC_KERNELS
> +                    || code == OACC_SERIAL)
> +                && OMP_CLAUSE_MAP_POINTS_TO_READONLY (c))
> +         {
> +           if (ctx->pt_readonly_ptrs == NULL)
> +             ctx->pt_readonly_ptrs = new hash_set<tree_operand_hash> ();
> +           ctx->pt_readonly_ptrs->add (OMP_CLAUSE_DECL (c));
> +         }
>         if (notice_outer)
>           goto do_notice;
>         break;

Also need to 'delete ctx->pt_readonly_ptrs;' somewhere.

> --- a/gcc/omp-low.cc
> +++ b/gcc/omp-low.cc
> @@ -14098,6 +14098,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
>               if (ref_to_array)
>                 x = fold_convert_loc (clause_loc, TREE_TYPE (new_var), x);
>               gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
> +             if (OMP_CLAUSE_MAP_POINTS_TO_READONLY (c) && VAR_P (x))
> +               DECL_POINTS_TO_READONLY (x) = 1;
>               if ((is_ref && !ref_to_array)
>                   || ref_to_ptr)
>                 {

This is in the middle of the
"Handle GOMP_MAP_FIRSTPRIVATE_{POINTER,REFERENCE} in second pass" code
block.  Again, for my understanding, is this empirically necessary, or a
systematic review?

> --- a/gcc/testsuite/c-c++-common/goacc/readonly-1.c
> +++ b/gcc/testsuite/c-c++-common/goacc/readonly-1.c
> @@ -19,8 +19,8 @@ int main (void)
>    return 0;
>  }
>
> -/* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*s.ptr \\\[len: 64\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: 128\\\]\\)" 1 "original" { target { c } } } } */
> -/* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: 64\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: 128\\\]\\)" 1 "original" { target { c++ } } } } */
> +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*s.ptr \\\[len: 64\\\]\\) map\\(pt_readonly,attach_detach:s.ptr \\\[bias: 0\\\]\\) map\\(readonly,to:x\\\[0\\\] \\\[len: 128\\\]\\) map\\(pt_readonly,firstprivate:x \\\[pointer assign, bias: 0\\\]\\)" 1 "original" { target { c } } } } */
> +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: 64\\\]\\) map\\(pt_readonly,attach_detach:s.ptr \\\[bias: 0\\\]\\) map\\(readonly,to:x\\\[0\\\] \\\[len: 128\\\]\\) map\\(pt_readonly,firstprivate:x \\\[pointer assign, bias: 0\\\]\\)" 1 "original" { target { c++ } } } } */
>  /* { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(readonly:x\\\[0\\\] \\\[len: 128\\\]\\);$" 1 "original" } } */

I suppose the new 'map(pt_readonly,attach_detach:s.ptr [bias: 0])' clause
was previously "hidden" in '.+'?  Please then change that in the first
patch "[PATCH, OpenACC 2.7] readonly modifier support in front-ends", so
that we can see here what actually is changing (only 'pt_readonly', I
suppose).

> --- /dev/null
> +++ b/gcc/testsuite/c-c++-common/goacc/readonly-2.c
> @@ -0,0 +1,15 @@
> +/* { dg-additional-options "-O -fdump-tree-fre" } */
> +
> +#pragma acc routine
> +extern void foo (int *ptr, int val);
> +
> +int main (void)
> +{
> +  int r, a[32];
> +  #pragma acc parallel copyin(readonly: a[:32]) copyout(r)
> +  {
> +    foo (a, a[8]);
> +    r = a[8];
> +  }
> +}
> +/* { dg-final { scan-tree-dump-times "r\.\[_0-9\]+ = MEM\\\[\[^_\]+_\[0-9\]+\\(ptro\\)\\\]\\\[8\\\];" 1 "fre1" } } */

Please add a comment why 'fre1', and what generally is being checked
here; that's not obvious to the casual reader.  (That is, me in a few
weeks.)  ;-)

Also add a scan for "before the optimization": two 'MEM's, I suppose?

> --- a/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90
> +++ b/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90
> @@ -20,8 +20,8 @@ program main
>    !$acc end parallel
>  end program main
>
> -! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\) .+ map\\(readonly,to:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\)" 1 "original" } }
> -! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:a\\\[\\(\\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) &a\\) / 4\\\] \\\[len: .+\\\]\\) .+ map\\(readonly,to:b\\\[\\(\\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) &b\\) / 4\\\] \\\[len: .+\\\]\\)" 1 "original" } }
> +! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\) map\\(pt_readonly,alloc:a.0 \\\[pointer assign, bias: \\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) a.0\\\]\\) map\\(readonly,to:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\) map\\(pt_readonly,alloc:b \\\[pointer assign, bias: \\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) b\\\]\\)" 1 "original" } }
> +! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:a\\\[\\(\\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) &a\\) / 4\\\] \\\[len: .+\\\]\\) map\\(pt_readonly,alloc:a \\\[pointer assign, bias: \\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) &a\\\]\\) map\\(readonly,to:b\\\[\\(\\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) &b\\) / 4\\\] \\\[len: .+\\\]\\) map\\(pt_readonly,alloc:b \\\[pointer assign, bias: \\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) &b\\\]\\)" 1 "original" } }
>  ! { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(readonly:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\) \\(readonly:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\);" 2 "original" } }

Same comment as for 'c-c++-common/goacc/readonly-1.c'.

> --- a/gcc/tree-pretty-print.cc
> +++ b/gcc/tree-pretty-print.cc
> @@ -907,6 +907,8 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
>        pp_string (pp, "map(");
>        if (OMP_CLAUSE_MAP_READONLY (clause))
>       pp_string (pp, "readonly,");
> +      if (OMP_CLAUSE_MAP_POINTS_TO_READONLY (clause))
> +     pp_string (pp, "pt_readonly,");
>        switch (OMP_CLAUSE_MAP_KIND (clause))
>       {
>       case GOMP_MAP_ALLOC:
> @@ -3436,6 +3438,8 @@ dump_generic_node (pretty_printer *pp, tree node, int spc, dump_flags_t flags,
>       pp_string (pp, "(D)");
>        if (SSA_NAME_OCCURS_IN_ABNORMAL_PHI (node))
>       pp_string (pp, "(ab)");
> +      if (SSA_NAME_POINTS_TO_READONLY_MEMORY (node))
> +     pp_string (pp, "(ptro)");
>        break;
>
>      case WITH_SIZE_EXPR:

> --- a/gcc/tree-ssanames.cc
> +++ b/gcc/tree-ssanames.cc
> @@ -402,6 +402,9 @@ make_ssa_name_fn (struct function *fn, tree var, gimple *stmt,
>    else
>      SSA_NAME_RANGE_INFO (t) = NULL;
>
> +  if (VAR_P (var) && DECL_POINTS_TO_READONLY (var))
> +    SSA_NAME_POINTS_TO_READONLY_MEMORY (t) = 1;
> +
>    SSA_NAME_IN_FREE_LIST (t) = 0;
>    SSA_NAME_IS_DEFAULT_DEF (t) = 0;
>    init_ssa_name_imm_use (t);

> --- a/gcc/tree.h
> +++ b/gcc/tree.h
> @@ -1021,6 +1021,13 @@ extern void omp_clause_range_check_failed (const_tree, const char *, int,
>  #define DECL_HIDDEN_STRING_LENGTH(NODE) \
>    (TREE_CHECK (NODE, PARM_DECL)->decl_common.decl_nonshareable_flag)
>
> +/* In a VAR_DECL, set for variables regarded as pointing to memory not written
> +   to. SSA_NAME_POINTS_TO_READONLY_MEMORY gets set for SSA_NAMEs created from
> +   such VAR_DECLs. Currently used by OpenACC 'readonly' modifier in copyin
> +   clauses.  */
> +#define DECL_POINTS_TO_READONLY(NODE) \
> +  (TREE_CHECK (NODE, VAR_DECL)->decl_common.decl_not_flexarray)
> +
>  /* In a CALL_EXPR, means that the call is the jump from a thunk to the
>     thunked-to function.  Be careful to avoid using this macro when one of the
>     next two applies instead.  */
> @@ -1815,6 +1822,10 @@ class auto_suppress_location_wrappers
>  #define OMP_CLAUSE_MAP_READONLY(NODE) \
>    TREE_READONLY (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
>
> +/* Set if 'OMP_CLAUSE_DECL (NODE)' points to read-only memory.  */
> +#define OMP_CLAUSE_MAP_POINTS_TO_READONLY(NODE) \
> +  TREE_CONSTANT (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_))

As in my "[PATCH, OpenACC 2.7] readonly modifier support in front-ends"
review, please document how certain flags are used for OMP clauses.


I note you're not actually using 'OMP_CLAUSE__CACHE__READONLY' anywhere
-- but that's OK given the current 'gcc/gimplify.cc:gimplify_oacc_cache'.
;-)


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

Re: [PATCH, OpenACC 2.7] Connect readonly modifier to points-to analysis

[Richard Biener]

On Fri, Oct 27, 2023 at 4:28 PM Thomas Schwinge <thomas@codesourcery.com> wrote:
>
> Hi!
>
> Richard, as the original author of 'SSA_NAME_POINTS_TO_READONLY_MEMORY':
> 2018 commit 6214d5c7e7470bdd5ecbeae668c2522551bfebbc (Subversion r263958)
> "Move const_parm trick to generic code"; 'gcc/tree.h':
>
>     /* Nonzero if this SSA_NAME is known to point to memory that may not
>        be written to.  This is set for default defs of function parameters
>        that have a corresponding r or R specification in the functions
>        fn spec attribute.  This is used by alias analysis.  */
>     #define SSA_NAME_POINTS_TO_READONLY_MEMORY(NODE) \
>         SSA_NAME_CHECK (NODE)->base.deprecated_flag
>
> ..., may I ask you to please help review the following patch
> (full-quoted)?
>
> For context: this patch here ("second patch") depends on a first patch:
> <inbox.sourceware.org/d0e6013f-ca38-b98d-dc01-b30adbd5901a@siemens.com>
> "[PATCH, OpenACC 2.7] readonly modifier support in front-ends".  That one
> is still under review/rework; so you're not able to apply this second
> patch here.
>
> In a nutshell: a 'readonly' modifier has been added to the OpenACC
> 'copyin' clause (copy host to device memory, don't copy back at end of
> region):
>
> | If the optional 'readonly' modifier appears, then the implementation may assume that the data
> | referenced by _var-list_ is never written to within the applicable region.
>
> That is, for example (untested):
>
>     #pragma acc routine
>     void escape(int *);
>
>     int x[32] = [...];
>     #pragma acc parallel copyin(readonly: x)
>     {
>       int a1 = x[3];
>       escape(x);
>       int a2 = x[3]; // Per 'readonly', don't need to reload 'x[3]' here.
>       //x[22] = 0; // Invalid -- but no diagnostic mandated.
>     }
>
> What Chung-Lin's first patch does is mark the OMP clause for 'x' (not the
> 'x' decl itself!) as 'readonly', via a new 'OMP_CLAUSE_MAP_READONLY'
> flag.
>
> The actual optimization then is done in this second patch.  Chung-Lin
> found that he could use 'SSA_NAME_POINTS_TO_READONLY_MEMORY' for that.
> I don't have much experience with most of the following generic code, so
> would appreciate a helping hand, whether that conceptually makes sense as
> well as from the implementation point of view:

No, I don't think you can use that flag on non-default-defs, nor
preserve it on copying.  So
it also doesn't nicely extend to DECLs as done by the patch.  We
currently _only_ use it
for incoming parameters.  When used on arbitrary code you can get to for example

ptr1(points-to-readony-memory) = &p->x;
... access via ptr1 ...
ptr2 = &p->x;
... access via ptr2 ...

where both are your OMP regions differently constrained (the constrain is on the
code in the region, _not_ on the actual protections of the pointed to
data, much like
for the fortran case).  But now CSE comes along and happily replaces all ptr2
with ptr2 in the second region and ... oops!

So no, re-using SSA_NAME_POINTS_TO_READONLY_MEMORY doesn't look good.

Richard.

> On 2023-07-25T23:52:06+0800, Chung-Lin Tang via Gcc-patches <gcc-patches@gcc.gnu.org> wrote:
> > On 2023/7/11 2:33 AM, Chung-Lin Tang via Gcc-patches wrote:
> >> As we discussed earlier, the work for actually linking this to middle-end
> >> points-to analysis is a somewhat non-trivial issue. This first patch allows
> >> the language feature to be used in OpenACC directives first (with no effect for now).
> >> The middle-end changes are probably going to be a later patch.
> >
> > This second patch tries to link the readonly modifier to points-to analysis.
> >
> > There already exists SSA_NAME_POINTS_TO_READONLY_MEMORY and it's support in the
> > alias oracle routines in tree-ssa-alias.cc, so basically what this patch does is
> > try to make the variables holding the array section base pointers to have this
> > flag set.
> >
> > There is an another OMP_CLAUSE_MAP_POINTS_TO_READONLY set by front-ends on the
> > associated pointer clauses if OMP_CLAUSE_MAP_READONLY is set.
> > Also a DECL_POINTS_TO_READONLY flag is set for VAR_DECLs when creating the tmp
> > vars carrying these receiver references on the offloaded side. These
> > eventually get translated to SSA_NAME_POINTS_TO_READONLY_MEMORY.
>
>
> > This still doesn't always work as expected in terms of optimization:
> > struct pointer fields and Fortran arrays (kind of like C structs) which have
> > several accesses to create the pointer access on the receive/offloaded side,
> > and SRA appears to not work on these sequences, so gets in the way of much
> > redundancy elimination.
>
> I understand correctly that this is left as future work?  Please add the test
> cases you have, XFAILed in some reasonable way.
>
>
> > Currently have one testcase where we can demonstrate 'readonly' can avoid
> > a clobber by function call.
>
> :-)
>
>
> > --- a/gcc/c/c-typeck.cc
> > +++ b/gcc/c/c-typeck.cc
> > @@ -14258,6 +14258,8 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
> >       OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
> >        else
> >       OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
> > +      if (OMP_CLAUSE_MAP_READONLY (c))
> > +     OMP_CLAUSE_MAP_POINTS_TO_READONLY (c2) = 1;
> >        OMP_CLAUSE_MAP_IMPLICIT (c2) = OMP_CLAUSE_MAP_IMPLICIT (c);
> >        if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
> >         && !c_mark_addressable (t))
>
> > --- a/gcc/cp/semantics.cc
> > +++ b/gcc/cp/semantics.cc
> > @@ -5872,6 +5872,8 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
> >           }
> >         else
> >           OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
> > +       if (OMP_CLAUSE_MAP_READONLY (c))
> > +         OMP_CLAUSE_MAP_POINTS_TO_READONLY (c2) = 1;
> >         OMP_CLAUSE_MAP_IMPLICIT (c2) = OMP_CLAUSE_MAP_IMPLICIT (c);
> >         if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
> >             && !cxx_mark_addressable (t))
>
> > --- a/gcc/fortran/trans-openmp.cc
> > +++ b/gcc/fortran/trans-openmp.cc
> > @@ -2524,6 +2524,8 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_exec_op op,
> >        node3 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
> >        OMP_CLAUSE_SET_MAP_KIND (node3, ptr_kind);
> >        OMP_CLAUSE_DECL (node3) = gfc_conv_descriptor_data_get (decl);
> > +      if (n->u.readonly)
> > +     OMP_CLAUSE_MAP_POINTS_TO_READONLY (node3) = 1;
> >        /* This purposely does not include GOMP_MAP_ALWAYS_POINTER.  The extra
> >        cast prevents gimplify.cc from recognising it as being part of the
> >        struct - and adding an 'alloc: for the 'desc.data' pointer, which
> > @@ -2559,6 +2561,8 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_exec_op op,
> >                               OMP_CLAUSE_MAP);
> >        OMP_CLAUSE_SET_MAP_KIND (node3, ptr_kind);
> >        OMP_CLAUSE_DECL (node3) = decl;
> > +      if (n->u.readonly)
> > +     OMP_CLAUSE_MAP_POINTS_TO_READONLY (node3) = 1;
> >      }
>
> Could combine these two into one, after
> 'if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))' reconverges here, like
> where 'OMP_CLAUSE_SIZE (node3)' is set:
>
> >    ptr2 = fold_convert (ptrdiff_type_node, ptr2);
> >    OMP_CLAUSE_SIZE (node3) = fold_build2 (MINUS_EXPR, ptrdiff_type_node,
>
> Is 'n->u.readonly == OMP_CLAUSE_MAP_READONLY (node)'?  If yes, would the
> latter be clearer to use as the 'if' expression (like in C, C++ front
> ends)?
>
> I see further additional 'OMP_CLAUSE_MAP' clauses synthesized, for
> example in 'gcc/cp/semantics.cc:handle_omp_array_sections', or
> 'gcc/fortran/trans-openmp.cc:gfc_trans_omp_array_section', also
> 'gcc/gimplify.cc'.  I assume these are not relevant to have
> 'OMP_CLAUSE_MAP_READONLY' -> 'OMP_CLAUSE_MAP_POINTS_TO_READONLY'
> propagated?  Actually, per your changes (see below), there is one
> 'OMP_CLAUSE_MAP_POINTS_TO_READONLY' propagation in
> 'gcc/gimplify.cc:build_omp_struct_comp_nodes'.
>
> Is the current situation re flag setting/propagation what was empirically
> necessary to make the test case work, or is it a systematic review?  (The
> former is fine; I'd just like to know.)
>
> > --- a/gcc/gimple-expr.cc
> > +++ b/gcc/gimple-expr.cc
> > @@ -376,6 +376,8 @@ copy_var_decl (tree var, tree name, tree type)
> >    DECL_CONTEXT (copy) = DECL_CONTEXT (var);
> >    TREE_USED (copy) = 1;
> >    DECL_SEEN_IN_BIND_EXPR_P (copy) = 1;
> > +  if (VAR_P (var))
> > +    DECL_POINTS_TO_READONLY (copy) = DECL_POINTS_TO_READONLY (var);
> >    DECL_ATTRIBUTES (copy) = DECL_ATTRIBUTES (var);
> >    if (DECL_USER_ALIGN (var))
> >      {
>
> > --- a/gcc/gimplify.cc
> > +++ b/gcc/gimplify.cc
> > @@ -221,6 +221,7 @@ struct gimplify_omp_ctx
> >    splay_tree variables;
> >    hash_set<tree> *privatized_types;
> >    tree clauses;
> > +  hash_set<tree_operand_hash> *pt_readonly_ptrs;
> >    /* Iteration variables in an OMP_FOR.  */
> >    vec<tree> loop_iter_var;
> >    location_t location;
> > @@ -628,6 +629,15 @@ internal_get_tmp_var (tree val, gimple_seq *pre_p, gimple_seq *post_p,
> >    gimplify_expr (&val, pre_p, post_p, is_gimple_reg_rhs_or_call,
> >                fb_rvalue);
> >
> > +  bool pt_readonly = false;
> > +  if (gimplify_omp_ctxp && gimplify_omp_ctxp->pt_readonly_ptrs)
> > +    {
> > +      tree ptr = val;
> > +      if (TREE_CODE (ptr) == POINTER_PLUS_EXPR)
> > +     ptr = TREE_OPERAND (ptr, 0);
> > +      pt_readonly = gimplify_omp_ctxp->pt_readonly_ptrs->contains (ptr);
> > +    }
>
> 'POINTER_PLUS_EXPR' is the only special thing we may run into, here?
> (Generally, I prefer 'if', 'else if, [...], 'else gcc_unreachable ()'.)
>
> > +
> >    if (allow_ssa
> >        && gimplify_ctxp->into_ssa
> >        && is_gimple_reg_type (TREE_TYPE (val)))
> > @@ -639,9 +649,18 @@ internal_get_tmp_var (tree val, gimple_seq *pre_p, gimple_seq *post_p,
> >         if (name)
> >           SET_SSA_NAME_VAR_OR_IDENTIFIER (t, create_tmp_var_name (name));
> >       }
> > +      if (pt_readonly)
> > +     SSA_NAME_POINTS_TO_READONLY_MEMORY (t) = 1;
> >      }
> >    else
> > -    t = lookup_tmp_var (val, is_formal, not_gimple_reg);
> > +    {
> > +      t = lookup_tmp_var (val, is_formal, not_gimple_reg);
> > +      if (pt_readonly)
> > +     {
> > +       DECL_POINTS_TO_READONLY (t) = 1;
> > +       gimplify_omp_ctxp->pt_readonly_ptrs->add (t);
> > +     }
> > +    }
> >
> >    mod = build2 (INIT_EXPR, TREE_TYPE (t), t, unshare_expr (val));
> >
> > @@ -8906,6 +8925,8 @@ build_omp_struct_comp_nodes (enum tree_code code, tree grp_start, tree grp_end,
> >    OMP_CLAUSE_SET_MAP_KIND (c2, mkind);
> >    OMP_CLAUSE_DECL (c2) = unshare_expr (OMP_CLAUSE_DECL (grp_end));
> >    OMP_CLAUSE_CHAIN (c2) = NULL_TREE;
> > +  if (OMP_CLAUSE_MAP_POINTS_TO_READONLY (grp_end))
> > +    OMP_CLAUSE_MAP_POINTS_TO_READONLY (c2) = 1;
> >    tree grp_mid = NULL_TREE;
> >    if (OMP_CLAUSE_CHAIN (grp_start) != grp_end)
> >      grp_mid = OMP_CLAUSE_CHAIN (grp_start);
>
> For my understanding, is this empirically necessary, or a systematic
> review?
>
> > @@ -11741,6 +11762,16 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
> >
> >             gimplify_omp_ctxp = outer_ctx;
> >           }
> > +       else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
> > +                && (code == OACC_PARALLEL
> > +                    || code == OACC_KERNELS
> > +                    || code == OACC_SERIAL)
> > +                && OMP_CLAUSE_MAP_POINTS_TO_READONLY (c))
> > +         {
> > +           if (ctx->pt_readonly_ptrs == NULL)
> > +             ctx->pt_readonly_ptrs = new hash_set<tree_operand_hash> ();
> > +           ctx->pt_readonly_ptrs->add (OMP_CLAUSE_DECL (c));
> > +         }
> >         if (notice_outer)
> >           goto do_notice;
> >         break;
>
> Also need to 'delete ctx->pt_readonly_ptrs;' somewhere.
>
> > --- a/gcc/omp-low.cc
> > +++ b/gcc/omp-low.cc
> > @@ -14098,6 +14098,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
> >               if (ref_to_array)
> >                 x = fold_convert_loc (clause_loc, TREE_TYPE (new_var), x);
> >               gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
> > +             if (OMP_CLAUSE_MAP_POINTS_TO_READONLY (c) && VAR_P (x))
> > +               DECL_POINTS_TO_READONLY (x) = 1;
> >               if ((is_ref && !ref_to_array)
> >                   || ref_to_ptr)
> >                 {
>
> This is in the middle of the
> "Handle GOMP_MAP_FIRSTPRIVATE_{POINTER,REFERENCE} in second pass" code
> block.  Again, for my understanding, is this empirically necessary, or a
> systematic review?
>
> > --- a/gcc/testsuite/c-c++-common/goacc/readonly-1.c
> > +++ b/gcc/testsuite/c-c++-common/goacc/readonly-1.c
> > @@ -19,8 +19,8 @@ int main (void)
> >    return 0;
> >  }
> >
> > -/* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*s.ptr \\\[len: 64\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: 128\\\]\\)" 1 "original" { target { c } } } } */
> > -/* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: 64\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: 128\\\]\\)" 1 "original" { target { c++ } } } } */
> > +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*s.ptr \\\[len: 64\\\]\\) map\\(pt_readonly,attach_detach:s.ptr \\\[bias: 0\\\]\\) map\\(readonly,to:x\\\[0\\\] \\\[len: 128\\\]\\) map\\(pt_readonly,firstprivate:x \\\[pointer assign, bias: 0\\\]\\)" 1 "original" { target { c } } } } */
> > +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: 64\\\]\\) map\\(pt_readonly,attach_detach:s.ptr \\\[bias: 0\\\]\\) map\\(readonly,to:x\\\[0\\\] \\\[len: 128\\\]\\) map\\(pt_readonly,firstprivate:x \\\[pointer assign, bias: 0\\\]\\)" 1 "original" { target { c++ } } } } */
> >  /* { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(readonly:x\\\[0\\\] \\\[len: 128\\\]\\);$" 1 "original" } } */
>
> I suppose the new 'map(pt_readonly,attach_detach:s.ptr [bias: 0])' clause
> was previously "hidden" in '.+'?  Please then change that in the first
> patch "[PATCH, OpenACC 2.7] readonly modifier support in front-ends", so
> that we can see here what actually is changing (only 'pt_readonly', I
> suppose).
>
> > --- /dev/null
> > +++ b/gcc/testsuite/c-c++-common/goacc/readonly-2.c
> > @@ -0,0 +1,15 @@
> > +/* { dg-additional-options "-O -fdump-tree-fre" } */
> > +
> > +#pragma acc routine
> > +extern void foo (int *ptr, int val);
> > +
> > +int main (void)
> > +{
> > +  int r, a[32];
> > +  #pragma acc parallel copyin(readonly: a[:32]) copyout(r)
> > +  {
> > +    foo (a, a[8]);
> > +    r = a[8];
> > +  }
> > +}
> > +/* { dg-final { scan-tree-dump-times "r\.\[_0-9\]+ = MEM\\\[\[^_\]+_\[0-9\]+\\(ptro\\)\\\]\\\[8\\\];" 1 "fre1" } } */
>
> Please add a comment why 'fre1', and what generally is being checked
> here; that's not obvious to the casual reader.  (That is, me in a few
> weeks.)  ;-)
>
> Also add a scan for "before the optimization": two 'MEM's, I suppose?
>
> > --- a/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90
> > +++ b/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90
> > @@ -20,8 +20,8 @@ program main
> >    !$acc end parallel
> >  end program main
> >
> > -! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\) .+ map\\(readonly,to:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\)" 1 "original" } }
> > -! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:a\\\[\\(\\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) &a\\) / 4\\\] \\\[len: .+\\\]\\) .+ map\\(readonly,to:b\\\[\\(\\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) &b\\) / 4\\\] \\\[len: .+\\\]\\)" 1 "original" } }
> > +! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\) map\\(pt_readonly,alloc:a.0 \\\[pointer assign, bias: \\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) a.0\\\]\\) map\\(readonly,to:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\) map\\(pt_readonly,alloc:b \\\[pointer assign, bias: \\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) b\\\]\\)" 1 "original" } }
> > +! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:a\\\[\\(\\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) &a\\) / 4\\\] \\\[len: .+\\\]\\) map\\(pt_readonly,alloc:a \\\[pointer assign, bias: \\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) &a\\\]\\) map\\(readonly,to:b\\\[\\(\\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) &b\\) / 4\\\] \\\[len: .+\\\]\\) map\\(pt_readonly,alloc:b \\\[pointer assign, bias: \\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) &b\\\]\\)" 1 "original" } }
> >  ! { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(readonly:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\) \\(readonly:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\);" 2 "original" } }
>
> Same comment as for 'c-c++-common/goacc/readonly-1.c'.
>
> > --- a/gcc/tree-pretty-print.cc
> > +++ b/gcc/tree-pretty-print.cc
> > @@ -907,6 +907,8 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
> >        pp_string (pp, "map(");
> >        if (OMP_CLAUSE_MAP_READONLY (clause))
> >       pp_string (pp, "readonly,");
> > +      if (OMP_CLAUSE_MAP_POINTS_TO_READONLY (clause))
> > +     pp_string (pp, "pt_readonly,");
> >        switch (OMP_CLAUSE_MAP_KIND (clause))
> >       {
> >       case GOMP_MAP_ALLOC:
> > @@ -3436,6 +3438,8 @@ dump_generic_node (pretty_printer *pp, tree node, int spc, dump_flags_t flags,
> >       pp_string (pp, "(D)");
> >        if (SSA_NAME_OCCURS_IN_ABNORMAL_PHI (node))
> >       pp_string (pp, "(ab)");
> > +      if (SSA_NAME_POINTS_TO_READONLY_MEMORY (node))
> > +     pp_string (pp, "(ptro)");
> >        break;
> >
> >      case WITH_SIZE_EXPR:
>
> > --- a/gcc/tree-ssanames.cc
> > +++ b/gcc/tree-ssanames.cc
> > @@ -402,6 +402,9 @@ make_ssa_name_fn (struct function *fn, tree var, gimple *stmt,
> >    else
> >      SSA_NAME_RANGE_INFO (t) = NULL;
> >
> > +  if (VAR_P (var) && DECL_POINTS_TO_READONLY (var))
> > +    SSA_NAME_POINTS_TO_READONLY_MEMORY (t) = 1;
> > +
> >    SSA_NAME_IN_FREE_LIST (t) = 0;
> >    SSA_NAME_IS_DEFAULT_DEF (t) = 0;
> >    init_ssa_name_imm_use (t);
>
> > --- a/gcc/tree.h
> > +++ b/gcc/tree.h
> > @@ -1021,6 +1021,13 @@ extern void omp_clause_range_check_failed (const_tree, const char *, int,
> >  #define DECL_HIDDEN_STRING_LENGTH(NODE) \
> >    (TREE_CHECK (NODE, PARM_DECL)->decl_common.decl_nonshareable_flag)
> >
> > +/* In a VAR_DECL, set for variables regarded as pointing to memory not written
> > +   to. SSA_NAME_POINTS_TO_READONLY_MEMORY gets set for SSA_NAMEs created from
> > +   such VAR_DECLs. Currently used by OpenACC 'readonly' modifier in copyin
> > +   clauses.  */
> > +#define DECL_POINTS_TO_READONLY(NODE) \
> > +  (TREE_CHECK (NODE, VAR_DECL)->decl_common.decl_not_flexarray)
> > +
> >  /* In a CALL_EXPR, means that the call is the jump from a thunk to the
> >     thunked-to function.  Be careful to avoid using this macro when one of the
> >     next two applies instead.  */
> > @@ -1815,6 +1822,10 @@ class auto_suppress_location_wrappers
> >  #define OMP_CLAUSE_MAP_READONLY(NODE) \
> >    TREE_READONLY (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
> >
> > +/* Set if 'OMP_CLAUSE_DECL (NODE)' points to read-only memory.  */
> > +#define OMP_CLAUSE_MAP_POINTS_TO_READONLY(NODE) \
> > +  TREE_CONSTANT (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_))
>
> As in my "[PATCH, OpenACC 2.7] readonly modifier support in front-ends"
> review, please document how certain flags are used for OMP clauses.
>
>
> I note you're not actually using 'OMP_CLAUSE__CACHE__READONLY' anywhere
> -- but that's OK given the current 'gcc/gimplify.cc:gimplify_oacc_cache'.
> ;-)
>
>
> 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