[OpenACC,2.7] readonly modifier support in front-ends
Checks
Commit Message
Hi Thomas,
this patch contains support for the 'readonly' modifier in copyin clauses
and the cache directive.
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.
(Also CCing Tobias because of the Fortran bits)
Tested on powerpc64le-linux with nvptx offloading. Is this okay for trunk?
Thanks,
Chung-Lin
2023-07-10 Chung-Lin Tang <cltang@codesourcery.com>
gcc/c/ChangeLog:
* c-parser.cc (c_parser_omp_var_list_parens):
Add 'bool *readonly = NULL' parameter, add readonly modifier parsing
support.
(c_parser_oacc_data_clause): Adjust c_parser_omp_var_list_parens call
to turn on readonly modifier parsing for copyin clause, set
OMP_CLAUSE_MAP_READONLY if readonly modifier found, update comments.
(c_parser_oacc_cache): Adjust c_parser_omp_var_list_parens call
to turn on readonly modifier parsing, set OMP_CLAUSE__CACHE__READONLY
if readonly modifier found, update comments.
gcc/cp/ChangeLog:
* parser.cc (cp_parser_omp_var_list):
Add 'bool *readonly = NULL' parameter, add readonly modifier parsing
support.
(cp_parser_oacc_data_clause): Adjust cp_parser_omp_var_list call
to turn on readonly modifier parsing for copyin clause, set
OMP_CLAUSE_MAP_READONLY if readonly modifier found, update comments.
(cp_parser_oacc_cache): Adjust cp_parser_omp_var_list call
to turn on readonly modifier parsing, set OMP_CLAUSE__CACHE__READONLY
if readonly modifier found, update comments.
gcc/fortran/ChangeLog:
* gfortran.h (typedef struct gfc_omp_namelist): Adjust map_op as
ENUM_BITFIELD field, add 'bool readonly' field.
* openmp.cc (gfc_match_omp_map_clause): Add 'bool readonly = false'
parameter, set n->u.readonly field.
(gfc_match_omp_clauses): Add readonly modifier parsing for OpenACC
copyin clause, adjust call to gfc_match_omp_map_clause.
(gfc_match_oacc_cache): Add readonly modifier parsing for OpenACC
cache directive, adjust call to gfc_match_omp_map_clause.
* trans-openmp.cc (gfc_trans_omp_clauses): Set OMP_CLAUSE_MAP_READONLY,
OMP_CLAUSE__CACHE__READONLY to 1 when readonly is set.
gcc/ChangeLog:
* tree-pretty-print.cc (dump_omp_clause): Add support for printing
OMP_CLAUSE_MAP_READONLY and OMP_CLAUSE__CACHE__READONLY.
* tree.h (OMP_CLAUSE_MAP_READONLY): New macro.
(OMP_CLAUSE__CACHE__READONLY): New macro.
gcc/testsuite/ChangeLog:
* c-c++-common/goacc/readonly-1.c: New test.
* gfortran.dg/goacc/readonly-1.f90: New test.
Comments
Hi,
just a remark regarding OpenMP. With
...omp ... firstprivate(var) allocator(omp_const_mem_alloc: var) one can also create constant memory in OpenMP.
Likewise with a custom allocator that uses the memory space
omp_const_mem_space, which is then a run-time thing. I don't think
that's particular useful on the host as the !PROT_WRITE property is a
memory-page thing which requires to allocate a multiple of a page size
(and after writing the value, mprotect can make it read only). But I
think it can be useful on the device (cf. OpenACC). OpenMP and OpenACC
likely differ in terms of whether an entry is in the mapping table
(firstprivate vs copy) and in the ref count. In any case, it would be
good to have the code written such that both OpenACC's and OpenMP's use
case can share as much code as possible, even if only OpenACC is
initially supported. Tobias PS: I should eventually have a closer look
at your patch!
On 10.07.23 20:33, Chung-Lin Tang wrote:
> this patch contains support for the 'readonly' modifier in copyin clauses
> and the cache directive.
>
> 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.
-----------------
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
Hi Chung-Lin, Tobias!
On 2023-07-11T02:33:58+0800, Chung-Lin Tang <chunglin.tang@siemens.com> wrote:
> this patch contains support for the 'readonly' modifier in copyin clauses
> and the cache directive.
Thanks!
> 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.
ACK.
> (Also CCing Tobias because of the Fortran bits)
A few specific GCC/Fortran questions for Tobias below, and some more
review comments for Chung-Lin:
> --- a/gcc/c/c-parser.cc
> +++ 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)
> {
> /* The clauses location. */
> location_t loc = c_parser_peek_token (parser)->location;
> @@ -14067,6 +14068,20 @@ c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
> matching_parens parens;
> if (parens.require_open (parser))
> {
> + if (readonly != NULL)
> + {
> + 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;
> + }
> + else
> + *readonly = false;
> + }
> list = c_parser_omp_variable_list (parser, loc, kind, list, allow_deref);
> parens.skip_until_found_close (parser);
> }
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. For example (random), like 'ancestor :' is parsed in
'c_parser_omp_clause_device', or 'conditional :' is parsed in
'c_parser_omp_clause_lastprivate'. (Yes, this does duplicate a bit of
code, but that's easy enough to follow along.)
The existing 'enum omp_clause_code kind', 'bool allow_deref' actually
affect the parsing process; the new 'bool readonly' only propagates a
flag.
> @@ -14084,7 +14099,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 +14154,22 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
> default:
> gcc_unreachable ();
> }
> +
> + /* Turn on readonly modifier parsing for copyin clause. */
> + bool readonly = false, *readonly_ptr = NULL;
> + if (c_kind == PRAGMA_OACC_CLAUSE_COPYIN)
> + readonly_ptr = &readonly;
> +
> tree nl, c;
> - nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list, true);
> + nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list, true,
> + readonly_ptr);
That is, similar to 'c_parser_omp_clause_device', or
'c_parser_omp_clause_lastprivate', inline 'c_parser_omp_var_list_parens'
here, and only for 'PRAGMA_OACC_CLAUSE_COPYIN' parse 'readonly :', then
(for all) use 'c_parser_omp_variable_list' etc. instead of
'c_parser_omp_var_list_parens', then set 'readonly':
> for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
> - OMP_CLAUSE_SET_MAP_KIND (c, kind);
> + {
> + OMP_CLAUSE_SET_MAP_KIND (c, kind);
> + if (readonly)
> + OMP_CLAUSE_MAP_READONLY (c) = 1;
> + }
>
> return nl;
> @@ -18212,6 +18242,9 @@ 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.
> */
>
> @@ -18219,8 +18252,14 @@ static tree
> c_parser_oacc_cache (location_t loc, c_parser *parser)
> {
> tree stmt, clauses;
> + bool readonly;
> +
> + clauses = c_parser_omp_var_list_parens (parser, OMP_CLAUSE__CACHE_, NULL,
> + false, &readonly);
> + 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);
Similarly.
> --- a/gcc/cp/parser.cc
> +++ b/gcc/cp/parser.cc
Similarly.
> --- a/gcc/fortran/gfortran.h
> +++ 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
> {
I did wonder whether the 'readonly' flag should live in the
'gfc_omp_namelist' (as done here -- similar to 'lastprivate_conditional',
for example), or in 'gfc_omp_clauses' (similar to 'ancestor', for
example). Then I realized/remembered that 'gfc_omp_clauses' exists only
once per directive (which is sufficient for 'ancestor', for example, as
there may be only one OpenMP 'device' clause), whereas 'gfc_omp_namelist'
exists once per list item -- which is what we need for 'readonly'. Thus,
the above looks good to me.
> --- a/gcc/fortran/openmp.cc
> +++ b/gcc/fortran/openmp.cc
> @@ -1196,7 +1196,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,
> @@ -1205,7 +1205,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.readonly = readonly;
> + }
> return true;
> }
Similar to 'c_parser_omp_var_list_parens' above, the existing
'bool allow_common', 'bool allow_derived' actually affect the parsing
process; the new 'bool readonly' only propagates a flag. Which I
acknowledge the existing 'gfc_omp_map_op map_op' also only does, but that
one's applicable to a lot more instances than 'readonly'. So I again
wonder if we should keep the latter out of 'gfc_match_omp_map_clause',
and instead set the flag when parsing the 'copyin' clauses; again, for
example (random), like 'ancestor :', or 'conditional :' are parsed --
which you're mostly already doing:
> @@ -2079,11 +2082,16 @@ 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 = false;
> + if (gfc_match ("readonly : ") == MATCH_YES)
> + readonly = true;
> + if (gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
> + OMP_MAP_TO, true,
> + allow_derived, readonly))
> + continue;
> + }
> }
..., so you'd just set 'readonly' here, instead of having
'gfc_match_omp_map_clause' do that. Care has to be taken to only do that
for the current list items, which you'll need 'gfc_omp_namelist *head'
for, or similar. Hmm. Effectively inline 'gfc_match_omp_map_clause'
here, or do add the 'bool readonly' argument to the latter, or something
else?
Or, we could add a new 'gcc/fortran/gfortran.h:gfc_omp_map_op' item
'OMP_MAP_TO_READONLY', which eventually translates into 'OMP_MAP_TO' with
'readonly' set? Then we'd just here call the (unaltered)
'gfc_match_omp_map_clause', with
'readonly ? OMP_MAP_TO_READONLY : OMP_MAP_TO'? Per
'git grep --cached '[^G]OMP_MAP_TO[^F]' -- gcc/fortran/' not a lot of
places need adjusting for that (most of the 'gcc/fortran/openmp.cc' ones
are not applicable).
Tobias?
> @@ -4008,20 +4016,35 @@ 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_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?
> + 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.readonly = true;
This already looks like how I thought it should look like.
> --- a/gcc/fortran/trans-openmp.cc
> +++ b/gcc/fortran/trans-openmp.cc
> @@ -3067,6 +3067,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
> || (n->expr && gfc_expr_attr (n->expr).pointer)))
> always_modifier = true;
>
> + if (n->u.readonly)
> + OMP_CLAUSE_MAP_READONLY (node) = 1;
> +
> switch (n->u.map_op)
> {
> case OMP_MAP_ALLOC:
> @@ -3920,6 +3923,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.readonly)
> + OMP_CLAUSE__CACHE__READONLY (node) = 1;
> omp_clauses = gfc_trans_add_clause (node, omp_clauses);
> }
> break;
> --- /dev/null
> +++ b/gcc/testsuite/c-c++-common/goacc/readonly-1.c
> @@ -0,0 +1,27 @@
> +/* { dg-additional-options "-fdump-tree-original" } */
> +
> +struct S
> +{
> + int *ptr;
> + float f;
> +};
> +
> +
> +int main (void)
> +{
> + int x[32];
> + struct S s = {x, 0};
> +
> + #pragma acc parallel copyin(readonly: x[:32], s.ptr[:16])
> + {
> + #pragma acc cache (readonly: x[:32])
> + }
> + 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 cache \\(readonly:x\\\[0\\\] \\\[len: 128\\\]\\);$" 1 "original" } } */
Are 'len: 64' etc. also correct for targets where 'sizeof (int) != 4'?
Maybe just mask these out; they're not the important thing we're testing
here?
> --- /dev/null
> +++ b/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90
> @@ -0,0 +1,28 @@
> +! { dg-additional-options "-fdump-tree-original" }
> +
> +subroutine foo (a, n)
> + integer :: n, a(:)
> + integer :: i, b(n)
> + !$acc parallel copyin(readonly: a(:), b(:n))
> + do i = 1,32
> + !$acc cache (readonly: a(:), b(:n))
> + enddo
> + !$acc end parallel
> +end subroutine foo
> +
> +program main
> + integer :: i, n = 32, a(32)
> + integer :: b(32)
> + !$acc parallel copyin(readonly: a(:32), b(:n))
> + do i = 1,32
> + !$acc cache (readonly: a(:), b(:n))
> + enddo
> + !$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 cache \\(readonly:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\) \\(readonly:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\);" 2 "original" } }
You're scanning only one of the two 'cache' directives? If that's
intentional, please add a comment, why. If not, add the missing
scanning.
Given the peculiarities of the Fortran parsing, where first all
directive's clauses are collected and then translated en bloc, I
suggest to extent the 'copyin' test cases to have several 'copyin'
clauses, some with, some without 'readonly' modifier, so we make sure
that 'readonly' is set only for the appropriate ones.
Generally, in addition to just 'parallel' compute construct, please
spread this out a bit, to also cover 'kernels', 'serial' compute
constructs, and the 'data' construct.
Generally, please also add testing for the 'declare' directive with
'copyin' with 'readonly' modifier -- and implement handling in case
that's not implicitly covered? (..., but please don't let you be dragged
into a number of pre-existing issues with OpenACC 'declare' -- I hope the
'readonly' handling is straightforward to test for.)
Given that per the implementation in the front ends, the handling of
'readonly' obviously -- famous last words? ;-) -- is specific to
'copyin', it's probably OK to not have test cases to verify that the
'readonly' modifier is rejected for other data clauses?
> --- a/gcc/tree-pretty-print.cc
> +++ b/gcc/tree-pretty-print.cc
> @@ -905,6 +905,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:
> @@ -1075,6 +1077,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;
> --- a/gcc/tree.h
> +++ b/gcc/tree.h
> @@ -1813,6 +1813,14 @@ class auto_suppress_location_wrappers
> #define OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE(NODE) \
> (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.addressable_flag)
>
> +/* Nonzero if OpenACC 'readonly' modifier set, used for 'copyin'. */
> +#define OMP_CLAUSE_MAP_READONLY(NODE) \
> + TREE_READONLY (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
> +
> +/* Same as above, for use in OpenACC cache directives. */
> +#define OMP_CLAUSE__CACHE__READONLY(NODE) \
> + TREE_READONLY (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE__CACHE_))
I'm not sure if these special accessor functions are actually useful, or
we should just directly use 'TREE_READONLY' instead? We're only using
them in contexts where it's clear that the 'OMP_CLAUSE_SUBCODE_CHECK' is
satisfied, for example.
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".
Setting 'TREE_READONLY' of the 'OMP_CLAUSE_DECL' instead of the clause
itself isn't the right thing to do -- or is it, and might already
indicate to the middle end the desired semantics? But does it maybe
conflict with front end/language-level use of 'TREE_READONLY' for 'const'
etc. (I suppose), and thus diagnostics for mismatches? I mean:
int a;
#pragma acc parallel copyin(readonly: a)
{
int *b = &a;
... should still continue to work (valid as long as '*b' isn't written
to), so should not raise any
"warning: initialization discards ‘const’ qualifier from pointer target type"
diagnostics. But if that's not a problem (I don't know how
'TREE_READONLY' is used elsewhere), maybe that's something to give a
thought to?
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?
Just some quick thoughts, obviously without any detailed analysis. ;-)
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
Hi Thomas & Chung-Lin,
On 20.07.23 15:33, Thomas Schwinge wrote:
> On 2023-07-11T02:33:58+0800, Chung-Lin Tang
> <chunglin.tang@siemens.com> wrote:
>> +++ b/gcc/c/c-parser.cc
>> @@ -14059,7 +14059,8 @@ c_parser_omp_variable_list (c_parser *parser,
>>
>> static tree
>> c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
>> - tree list, bool allow_deref = false)
>> + tree list, bool allow_deref = false,
>> + bool *readonly = NULL)
>> ...
> Instead of doing this in 'c_parser_omp_var_list_parens', I think it's
> clearer to have this special 'readonly :' parsing logic in the two places
> where it's used.
I concur. The same issue also occurred for OpenMP's
c_parser_omp_clause_to, and c_parser_omp_clause_from and the 'present'
modifier. For it, I created a combined function but the main reason for
that is that OpenMP also permits more modifiers (like 'iterators'),
which would cause more duplication of code ('iterator' is not yet
supported).
For something as simple to parse as this modifier, I would just do it at
the two places – as Thomas suggested.
>> +++ b/gcc/fortran/gfortran.h
>> @@ -1360,7 +1360,11 @@ typedef struct gfc_omp_namelist
>> {
>> gfc_omp_reduction_op reduction_op;
>> gfc_omp_depend_doacross_op depend_doacross_op;
>> - gfc_omp_map_op map_op;
>> + struct
>> + {
>> + ENUM_BITFIELD (gfc_omp_map_op) map_op:8;
>> + bool readonly;
>> + };
>> gfc_expr *align;
>> struct
>> {
> [...] Thus, the above looks good to me.
I concur but I wonder whether it would be cleaner to name the struct;
this makes it also more obvious what belongs together in the union.
Namely, naming the struct 'map' and then changing the 45 users from
'u.map_op' to 'u.map.op' and the new 'u.readonly' to 'u.map.readonly'. –
this seems to be cleaner.
>> static bool
>> gfc_match_omp_map_clause (gfc_omp_namelist **list, gfc_omp_map_op map_op,
>> - bool allow_common, bool allow_derived)
>> + bool allow_common, bool allow_derived, bool readonly = false)
>> {
> Similar to 'c_parser_omp_var_list_parens' above,
I concur that not doing it here is cleaner.
> again, for
> example (random), like 'ancestor :', or 'conditional :' are parsed --
> which you're mostly already doing
I think OpenMP's "present" (as modifier to "omp target updates"'s
"to"/"from") is a better example than "ancestor" as for present we also
have a list. See: gfc_match_motion_var_list how to handle the headp.
(There an extra functions was used as in the future also other modifiers
like 'iterator' will be used.)
However, as Thomas noted, the patch contains also an example (see
further down in Thomas' email, not quoted here).
> Or, we could add a new 'gcc/fortran/gfortran.h:gfc_omp_map_op' item
> 'OMP_MAP_TO_READONLY', which eventually translates into 'OMP_MAP_TO' with
> 'readonly' set?
I think having the additional flag is easier to understand - and at least
memory wise we do not save memory as it is in a union. The advantage
of not having a union is that accessing the int-enum is faster than accessing
an char-wide bitset enum.
In terms of code changes (and without having a closer look), the two
approaches seems to be be similar.
Hence, using OMP_MAP_TO_READONLY for OpenACC would be fine, too. And
I do not have a strong preference for either.
* * *
I did wonder about the following, but I now believe it won't affect
the choice. Namely, we want to handle at some point the following:
!$omp target firstprivate(var) allocator(omp_const_mem_alloc: var)
This could be turned into GOMP_MAP_FIRSTPRIVATE... + OMP_.*READONLY flag.
But if we don't do it in the FE, the internal Fortran representation
does not matter.
Advantage for doing it in the ME: Only one code location, especially as
we might use the opportunity to also check that the omp_const_mem_alloc
is only used with privatization (in OpenMP).
Difference: OpenMP uses 'firstprivate' (i.e. private copy, no reference count bump,
only permitted for 'target') while OpenACC uses 'copy' which implies reference
counting and permitted in 'acc (enter/exit) data' and not only for compute constructs.
OpenMP in principle also permits user-defined allocator with a constant
memory space - I am not completely sure whether/when it can be used with
omp target firstprivate(...) allocator(my_alloc : ...)
> Then we'd just here call the (unaltered)
> 'gfc_match_omp_map_clause', with
> 'readonly ? OMP_MAP_TO_READONLY : OMP_MAP_TO'? Per
> 'git grep --cached '[^G]OMP_MAP_TO[^F]' -- gcc/fortran/' not a lot of
> places need adjusting for that (most of the 'gcc/fortran/openmp.cc' ones
> are not applicable).
I think either would work. – I have no strong feeling what's better.
But you still need to handle it for clause resolution.
> + if (gfc_match ("readonly :") == MATCH_YES)
> I note this one does not have a space after ':' in 'gfc_match', but the
> one above in 'gfc_match_omp_clauses' does. I don't know off-hand if that
> makes a difference in parsing -- probably not, as all of
> 'gcc/fortran/openmp.cc' generally doesn't seem to be very consistent
> about these two variants?
It *does* make a difference. And for obvious reasons. You don't want to permit:
!$acc kernels asnyccopy(a)
but require at least one space (or comma) between "async" and "copy"..
(In fixed form Fortran, it would be fine - as would be "!$acc k e nelsasy nc co p y(a)".)
A " " matches zero or more whitespaces, but with gfc_match_space you can find out
whether there was whitespace or not.
Whether the tailing " " in the gfc_match matters or not, depends on what comes next.
If there is a "gfc_gobble_whitespace ();", everything is fine. If not, the next to match
has to start with a " ", which is usually ugly; an exception is " , " or " )" which still
is somewhat fine.
I think that it is mostly implemented correctly, but I wouldn't be surprised if a
space is missing in some matches - be it a tailing white space or e.g. in "foo:" before
the colon.
BTW: One reason of stripping tailing spaces before matching a non-whitespace: the
associated location is the one before the parsing; thus, for a match error or when saving
the old_locus, pointing to the first non-whitespace looks nicer than pointing to the
(first of the) whitspace character(s).
>> + if (readonly)
>> + for (gfc_omp_namelist *n = *head; n; n = n->next)
>> + n->u.readonly = true;
> This already looks like how I thought it should look like.
Indeed.--- a/gcc/tree.h
>> +++ b/gcc/tree.h
>> @@ -1813,6 +1813,14 @@ class auto_suppress_location_wrappers
>> #define OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE(NODE) \
>> (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.addressable_flag)
>>
>> +/* Nonzero if OpenACC 'readonly' modifier set, used for 'copyin'. */
>> +#define OMP_CLAUSE_MAP_READONLY(NODE) \
>> + TREE_READONLY (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
>> +
>> +/* Same as above, for use in OpenACC cache directives. */
>> +#define OMP_CLAUSE__CACHE__READONLY(NODE) \
>> + TREE_READONLY (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE__CACHE_))
> I'm not sure if these special accessor functions are actually useful, or
> we should just directly use 'TREE_READONLY' instead? We're only using
> them in contexts where it's clear that the 'OMP_CLAUSE_SUBCODE_CHECK' is
> satisfied, for example.
I find directly using TREE_READONLY confusing.
> Setting 'TREE_READONLY' of the 'OMP_CLAUSE_DECL' instead of the clause
> itself isn't the right thing to do -- or is it, and might already
> indicate to the middle end the desired semantics? But does it maybe
> conflict with front end/language-level use of 'TREE_READONLY' for 'const'
> etc. (I suppose), and thus diagnostics for mismatches?
I think is is cleaner not to one flag to mean two different things.
In particular, wouldn't the following cause issues, if you mark 'a' as TREE_READONLY?
int a;
#pragma acc parallel copyin(readonly : a)
{...}
a = 5;
> Or, early in the middle end, propagate 'TREE_READONLY' from the clause to
> its 'OMP_CLAUSE_DECL'? Might need to 'unshare_expr' the latter for
> modification and use in the associated region only?
Unsharing a tree would surely help – but it is still ugly and, for
declarations, unshare_expr does not create a copy!
Tobias
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
@@ -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)
{
/* The clauses location. */
location_t loc = c_parser_peek_token (parser)->location;
@@ -14067,6 +14068,20 @@ c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
matching_parens parens;
if (parens.require_open (parser))
{
+ if (readonly != NULL)
+ {
+ 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;
+ }
+ else
+ *readonly = false;
+ }
list = c_parser_omp_variable_list (parser, loc, kind, list, allow_deref);
parens.skip_until_found_close (parser);
}
@@ -14084,7 +14099,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 +14154,22 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
default:
gcc_unreachable ();
}
+
+ /* Turn on readonly modifier parsing for copyin clause. */
+ bool readonly = false, *readonly_ptr = NULL;
+ if (c_kind == PRAGMA_OACC_CLAUSE_COPYIN)
+ readonly_ptr = &readonly;
+
tree nl, c;
- nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list, true);
+ nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list, true,
+ readonly_ptr);
for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
- OMP_CLAUSE_SET_MAP_KIND (c, kind);
+ {
+ OMP_CLAUSE_SET_MAP_KIND (c, kind);
+ if (readonly)
+ OMP_CLAUSE_MAP_READONLY (c) = 1;
+ }
return nl;
}
@@ -18212,6 +18242,9 @@ 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.
*/
@@ -18219,8 +18252,14 @@ static tree
c_parser_oacc_cache (location_t loc, c_parser *parser)
{
tree stmt, clauses;
+ bool readonly;
+
+ clauses = c_parser_omp_var_list_parens (parser, OMP_CLAUSE__CACHE_, NULL,
+ false, &readonly);
+ 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);
@@ -37727,11 +37727,27 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
static tree
cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list,
- bool allow_deref = false)
+ bool allow_deref = false, bool *readonly = NULL)
{
if (cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
- return cp_parser_omp_var_list_no_open (parser, kind, list, NULL,
- allow_deref);
+ {
+ if (readonly != NULL)
+ {
+ cp_token *token = cp_lexer_peek_token (parser->lexer);
+ if (token->type == CPP_NAME
+ && !strcmp (IDENTIFIER_POINTER (token->u.value), "readonly")
+ && cp_lexer_nth_token_is (parser->lexer, 2, CPP_COLON))
+ {
+ cp_lexer_consume_token (parser->lexer);
+ cp_lexer_consume_token (parser->lexer);
+ *readonly = true;
+ }
+ else
+ *readonly = false;
+ }
+ return cp_parser_omp_var_list_no_open (parser, kind, list, NULL,
+ allow_deref);
+ }
return list;
}
@@ -37746,7 +37762,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,
@@ -37797,11 +37817,22 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
default:
gcc_unreachable ();
}
+
+ /* Turn on readonly modifier parsing for copyin clause. */
+ bool readonly = false, *readonly_ptr = NULL;
+ if (c_kind == PRAGMA_OACC_CLAUSE_COPYIN)
+ readonly_ptr = &readonly;
+
tree nl, c;
- nl = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list, true);
+ nl = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list, true,
+ readonly_ptr);
for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
- OMP_CLAUSE_SET_MAP_KIND (c, kind);
+ {
+ OMP_CLAUSE_SET_MAP_KIND (c, kind);
+ if (readonly)
+ OMP_CLAUSE_MAP_READONLY (c) = 1;
+ }
return nl;
}
@@ -45875,6 +45906,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
@@ -45885,8 +45919,14 @@ cp_parser_oacc_cache (cp_parser *parser, cp_token *pragma_tok)
auto_suppress_location_wrappers sentinel;
tree stmt, clauses;
+ bool readonly;
+
+ clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE__CACHE_, NULL,
+ false, &readonly);
+ 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));
@@ -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
{
@@ -1196,7 +1196,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,
@@ -1205,7 +1205,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.readonly = readonly;
+ }
return true;
}
@@ -2079,11 +2082,16 @@ 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 = false;
+ if (gfc_match ("readonly : ") == MATCH_YES)
+ readonly = true;
+ if (gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
+ OMP_MAP_TO, true,
+ allow_derived, readonly))
+ continue;
+ }
}
else if (gfc_match_omp_variable_list ("copyin (",
&c->lists[OMP_LIST_COPYIN],
@@ -4008,20 +4016,35 @@ 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_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.readonly = true;
+
if (gfc_current_state() != COMP_DO
&& gfc_current_state() != COMP_DO_CONCURRENT)
{
@@ -3067,6 +3067,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
|| (n->expr && gfc_expr_attr (n->expr).pointer)))
always_modifier = true;
+ if (n->u.readonly)
+ OMP_CLAUSE_MAP_READONLY (node) = 1;
+
switch (n->u.map_op)
{
case OMP_MAP_ALLOC:
@@ -3920,6 +3923,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.readonly)
+ OMP_CLAUSE__CACHE__READONLY (node) = 1;
omp_clauses = gfc_trans_add_clause (node, omp_clauses);
}
break;
new file mode 100644
@@ -0,0 +1,27 @@
+/* { dg-additional-options "-fdump-tree-original" } */
+
+struct S
+{
+ int *ptr;
+ float f;
+};
+
+
+int main (void)
+{
+ int x[32];
+ struct S s = {x, 0};
+
+ #pragma acc parallel copyin(readonly: x[:32], s.ptr[:16])
+ {
+ #pragma acc cache (readonly: x[:32])
+ }
+ 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 cache \\(readonly:x\\\[0\\\] \\\[len: 128\\\]\\);$" 1 "original" } } */
+
+
+
new file mode 100644
@@ -0,0 +1,28 @@
+! { dg-additional-options "-fdump-tree-original" }
+
+subroutine foo (a, n)
+ integer :: n, a(:)
+ integer :: i, b(n)
+ !$acc parallel copyin(readonly: a(:), b(:n))
+ do i = 1,32
+ !$acc cache (readonly: a(:), b(:n))
+ enddo
+ !$acc end parallel
+end subroutine foo
+
+program main
+ integer :: i, n = 32, a(32)
+ integer :: b(32)
+ !$acc parallel copyin(readonly: a(:32), b(:n))
+ do i = 1,32
+ !$acc cache (readonly: a(:), b(:n))
+ enddo
+ !$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 cache \\(readonly:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\) \\(readonly:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\);" 2 "original" } }
+
+
+
@@ -905,6 +905,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:
@@ -1075,6 +1077,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;
@@ -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_))
+
/* True on an OMP_CLAUSE_USE_DEVICE_PTR with an OpenACC 'if_present'
clause. */
#define OMP_CLAUSE_USE_DEVICE_PTR_IF_PRESENT(NODE) \