From: Thomas Schwinge <tschwinge@baylibre.com>
To: Chung-Lin Tang <cltang@baylibre.com>
Cc: Tobias Burnus <tburnus@baylibre.com>,
gcc-patches@gcc.gnu.org, fortran@gcc.gnu.org
Subject: Re: [PATCH, OpenACC 2.7, v2] readonly modifier support in front-ends
Date: Wed, 13 Mar 2024 10:12:17 +0100 [thread overview]
Message-ID: <87le6mebri.fsf@euler.schwinge.ddns.net> (raw)
In-Reply-To: <2f568f7d-807b-41d1-befb-40039e2edb74@pllab.cs.nthu.edu.tw>
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) \
next prev parent reply other threads:[~2024-03-13 9:12 UTC|newest]
Thread overview: 8+ messages / expand[flat|nested] mbox.gz Atom feed top
[not found] <d0e6013f-ca38-b98d-dc01-b30adbd5901a@siemens.com>
[not found] ` <87lefaaesb.fsf@euler.schwinge.homeip.net>
[not found] ` <b45bd2bb-a306-5e42-7c84-53433e8d06a2@codesourcery.com>
[not found] ` <b5af4407-1538-802f-92ca-aae843258c15@siemens.com>
2023-10-26 9:43 ` Thomas Schwinge
2024-03-07 8:02 ` Chung-Lin Tang
2024-03-13 9:12 ` Thomas Schwinge [this message]
2024-03-14 15:09 ` 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
2024-03-14 16:55 ` OpenACC 2.7: front-end support for readonly modifier: Add basic OpenACC 'declare' testing Tobias Burnus
2024-03-14 16:55 ` Tobias Burnus
[not found] ` <5196826c-e81a-ab5c-63e9-bd8509232da0@siemens.com>
2023-10-27 14:28 ` [PATCH, OpenACC 2.7] Connect readonly modifier to points-to analysis Thomas Schwinge
2023-10-30 12:46 ` Richard Biener
Reply instructions:
You may reply publicly to this message via plain-text email
using any one of the following methods:
* Save the following mbox file, import it into your mail client,
and reply-to-all from there: mbox
Avoid top-posting and favor interleaved quoting:
https://en.wikipedia.org/wiki/Posting_style#Interleaved_style
* Reply using the --to, --cc, and --in-reply-to
switches of git-send-email(1):
git send-email \
--in-reply-to=87le6mebri.fsf@euler.schwinge.ddns.net \
--to=tschwinge@baylibre.com \
--cc=cltang@baylibre.com \
--cc=fortran@gcc.gnu.org \
--cc=gcc-patches@gcc.gnu.org \
--cc=tburnus@baylibre.com \
/path/to/YOUR_REPLY
https://kernel.org/pub/software/scm/git/docs/git-send-email.html
* If your mail client supports setting the In-Reply-To header
via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line
before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).