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
@@ -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;
+ 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;
+ 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);
@@ -37750,7 +37750,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,
@@ -37801,11 +37805,33 @@ 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, true);
+ 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,
+ true);
+ }
- for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
- OMP_CLAUSE_SET_MAP_KIND (c, kind);
+ 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;
}
@@ -45825,6 +45851,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
@@ -45834,9 +45863,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));
@@ -1398,6 +1398,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)
{
@@ -1465,7 +1468,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;
@@ -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) op:8;
+ bool readonly;
+ } map;
gfc_expr *align;
struct
{
@@ -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;
}
@@ -1520,7 +1523,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)
@@ -2180,11 +2183,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],
@@ -3101,7 +3109,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;
@@ -3942,7 +3950,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);
@@ -3976,7 +3984,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:
@@ -4091,20 +4099,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.map.readonly = true;
+
if (gfc_current_state() != COMP_DO
&& gfc_current_state() != COMP_DO_CONCURRENT)
{
@@ -8142,8 +8165,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)
@@ -8153,7 +8176,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
@@ -8224,7 +8247,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);
@@ -8246,7 +8269,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:
@@ -8273,7 +8296,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:
@@ -8283,16 +8306,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 "
@@ -8302,7 +8325,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:
@@ -8312,16 +8335,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 "
@@ -6614,7 +6614,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 ();
@@ -6716,10 +6716,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:
@@ -3067,7 +3067,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);
@@ -3194,8 +3197,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);
@@ -3261,7 +3264,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;
@@ -3284,7 +3287,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;
@@ -3316,18 +3319,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);
@@ -3345,7 +3348,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
@@ -3361,7 +3364,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
@@ -3553,8 +3556,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));
@@ -3585,7 +3588,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
fold_convert (size_type_node,
se.string_length),
TYPE_SIZE_UNIT (tmp));
- 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;
@@ -3642,8 +3645,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);
@@ -3689,8 +3692,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))))
{
@@ -3724,8 +3727,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)
map_kind = GOMP_MAP_RELEASE;
@@ -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.map.readonly)
+ OMP_CLAUSE__CACHE__READONLY (node) = 1;
omp_clauses = gfc_trans_add_clause (node, omp_clauses);
}
break;
@@ -6333,7 +6338,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;
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;
@@ -1811,6 +1811,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_NOTHROW (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
+
+/* Same as above, for use in OpenACC cache directives. */
+#define OMP_CLAUSE__CACHE__READONLY(NODE) \
+ TREE_NOTHROW (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) \