| Message ID | b5af4407-1538-802f-92ca-aae843258c15@siemens.com |
|---|---|
| State | New |
| Headers | show |
| Series | [OpenACC,2.7,v2] readonly modifier support in front-ends | expand |
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
Hi Thomas, Tobias, On 2023/10/26 6:43 PM, Thomas Schwinge wrote: >>>>> +++ b/gcc/tree.h >>>>> @@ -1813,6 +1813,14 @@ class auto_suppress_location_wrappers >>>>> #define OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE(NODE) \ >>>>> (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.addressable_flag) >>>>> >>>>> +/* Nonzero if OpenACC 'readonly' modifier set, used for 'copyin'. */ >>>>> +#define OMP_CLAUSE_MAP_READONLY(NODE) \ >>>>> + TREE_READONLY (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)) >>>>> + >>>>> +/* Same as above, for use in OpenACC cache directives. */ >>>>> +#define OMP_CLAUSE__CACHE__READONLY(NODE) \ >>>>> + TREE_READONLY (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE__CACHE_)) >>>> I'm not sure if these special accessor functions are actually useful, or >>>> we should just directly use 'TREE_READONLY' instead? We're only using >>>> them in contexts where it's clear that the 'OMP_CLAUSE_SUBCODE_CHECK' is >>>> satisfied, for example. >>> I find directly using TREE_READONLY confusing. >> >> FWIW, I've changed to use TREE_NOTHROW instead, if it can give a better sense of safety :P > > I don't understand that, why not use 'TREE_READONLY'? > >> I think there's a misunderstanding here anyways: we are not relying on a DECL marked >> TREE_READONLY here. We merely need the OMP_CLAUSE_MAP to be marked as OMP_CLAUSE_MAP_READONLY == 1. > > Yes, I understand that. My question was why we don't just use > 'TREE_READONLY (c)', where 'c' is the > 'OMP_CLAUSE_MAP'/'OMP_CLAUSE__CACHE_' clause (not its decl), and avoid > the indirection through > '#define OMP_CLAUSE_MAP_READONLY'/'#define OMP_CLAUSE__CACHE__READONLY', > given that we're only using them in contexts where it's clear that the > 'OMP_CLAUSE_SUBCODE_CHECK' is satisfied. I don't have a strong > preference, though. After further re-testing using TREE_NOTHROW, I have reverted to using TREE_READONLY, because TREE_NOTHROW clashes with OMP_CLAUSE_RELEASE_DESCRIPTOR (which doesn't use the OMP_CLAUSE_MAP_* naming convention and is not documented in gcc/tree-core.h either, hmmm...) I have added the comment adjustments in gcc/tree-core.h for the new uses of TREE_READONLY/readonly_flag. We basically all use OMP_CLAUSE_SUBCODE_CHECK macros for OpenMP clause expressions exclusively, so I don't see a reason to diverge from that style (even when context is clear). > Either way, you still need to document this: > > | Also, for the new use for OMP clauses, update 'gcc/tree.h:TREE_READONLY', > | and in 'gcc/tree-core.h' for 'readonly_flag' the > | "table lists the uses of each of the above flags". Okay, done as mentioned above. > In addition to a few individual comments above and below, you've also not > yet responded to my requests re test cases. I have greatly expanded the test scan patterns to include parallel/kernels/serial/data/enter data, as well as non-readonly copyin clause together with readonly. Also added simple 'declare' tests, but there is not anything to scan in the 'tree-original' dump though. >> + tree nl = list; >> + bool readonly = false; >> + matching_parens parens; >> + if (parens.require_open (parser)) >> + { >> + /* Turn on readonly modifier parsing for copyin clause. */ >> + if (c_kind == PRAGMA_OACC_CLAUSE_COPYIN) >> + { >> + c_token *token = c_parser_peek_token (parser); >> + if (token->type == CPP_NAME >> + && !strcmp (IDENTIFIER_POINTER (token->value), "readonly") >> + && c_parser_peek_2nd_token (parser)->type == CPP_COLON) >> + { >> + c_parser_consume_token (parser); >> + c_parser_consume_token (parser); >> + readonly = true; >> + } >> + } >> + location_t loc = c_parser_peek_token (parser)->location; > > I suppose 'loc' here now points to after the opening '(' or after the > 'readonly :'? This is different from what 'c_parser_omp_var_list_parens' > does, and indeed, 'c_parser_omp_variable_list' states that "CLAUSE_LOC is > the location of the clause", not the location of the variable-list? As > this, I suppose, may change diagnostics, please restore the original > behavior. (This appears to be different in the C++ front end, huh.) Thanks for catching this! Fixed. >> --- a/gcc/fortran/openmp.cc >> +++ b/gcc/fortran/openmp.cc >> @@ -1197,7 +1197,7 @@ omp_inv_mask::omp_inv_mask (const omp_mask &m) : omp_mask (m) >> >> static bool >> gfc_match_omp_map_clause (gfc_omp_namelist **list, gfc_omp_map_op map_op, >> - bool allow_common, bool allow_derived) >> + bool allow_common, bool allow_derived, bool readonly = false) >> { >> gfc_omp_namelist **head = NULL; >> if (gfc_match_omp_variable_list ("", list, allow_common, NULL, &head, true, >> @@ -1206,7 +1206,10 @@ gfc_match_omp_map_clause (gfc_omp_namelist **list, gfc_omp_map_op map_op, >> { >> gfc_omp_namelist *n; >> for (n = *head; n; n = n->next) >> - n->u.map_op = map_op; >> + { >> + n->u.map.op = map_op; >> + n->u.map.readonly = readonly; >> + } >> return true; >> } > > Didn't we conclude that "not doing it here is cleaner" (Tobias' words), > and instead do this "Similar to 'c_parser_omp_var_list_parens'" (my > words)? That is, not add the 'bool readonly' formal parameter to > 'gfc_match_omp_map_clause'. Fixed in this v3 patch. Again, tested on x86_64-linux + nvptx offloading. Okay for mainline? Thanks, Chung-Lin gcc/c/ChangeLog: * c-parser.cc (c_parser_oacc_data_clause): Add parsing support for 'readonly' modifier, set OMP_CLAUSE_MAP_READONLY if readonly modifier found, update comments. (c_parser_oacc_cache): Add parsing support for 'readonly' modifier, set OMP_CLAUSE__CACHE__READONLY if readonly modifier found, update comments. gcc/cp/ChangeLog: * parser.cc (cp_parser_oacc_data_clause): Add parsing support for 'readonly' modifier, set OMP_CLAUSE_MAP_READONLY if readonly modifier found, update comments. (cp_parser_oacc_cache): Add parsing support for 'readonly' modifier, set OMP_CLAUSE__CACHE__READONLY if readonly modifier found, update comments. gcc/fortran/ChangeLog: * dump-parse-tree.cc (show_omp_namelist): Print "readonly," for OMP_LIST_MAP and OMP_LIST_CACHE if n->u.map.readonly is set. Adjust 'n->u.map_op' to 'n->u.map.op'. * gfortran.h (typedef struct gfc_omp_namelist): Adjust map_op as 'ENUM_BITFIELD (gfc_omp_map_op) op:8', add 'bool readonly' field, change to named struct field 'map'. * openmp.cc (gfc_match_omp_map_clause): Adjust 'n->u.map_op' to 'n->u.map.op'. (gfc_match_omp_clause_reduction): Likewise. (gfc_match_omp_clauses): Add readonly modifier parsing for OpenACC copyin clause, set 'n->u.map.op' and 'n->u.map.readonly' for parsed clause. Adjust 'n->u.map_op' to 'n->u.map.op'. (gfc_match_oacc_declare): Adjust 'n->u.map_op' to 'n->u.map.op'. (gfc_match_oacc_cache): Add readonly modifier parsing for OpenACC cache directive. (resolve_omp_clauses): Adjust 'n->u.map_op' to 'n->u.map.op'. * trans-decl.cc (add_clause): Adjust 'n->u.map_op' to 'n->u.map.op'. (finish_oacc_declare): Likewise. * trans-openmp.cc (gfc_trans_omp_clauses): Set OMP_CLAUSE_MAP_READONLY, OMP_CLAUSE__CACHE__READONLY to 1 when readonly is set. Adjust 'n->u.map_op' to 'n->u.map.op'. (gfc_add_clause_implicitly): Adjust 'n->u.map_op' to 'n->u.map.op'. gcc/ChangeLog: * tree.h (OMP_CLAUSE_MAP_READONLY): New macro. (OMP_CLAUSE__CACHE__READONLY): New macro. * tree-core.h (struct GTY(()) tree_base): Adjust comments for new uses of readonly_flag bit in OMP_CLAUSE_MAP_READONLY and OMP_CLAUSE__CACHE__READONLY. * tree-pretty-print.cc (dump_omp_clause): Add support for printing OMP_CLAUSE_MAP_READONLY and OMP_CLAUSE__CACHE__READONLY. gcc/testsuite/ChangeLog: * c-c++-common/goacc/readonly-1.c: New test. * gfortran.dg/goacc/readonly-1.f90: New test. 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) \
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) \
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index 24a6eb6e459..5779f499ae1 100644 --- 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; + 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); diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index d7ef5b34d42..ac8a656874a 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -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)); diff --git a/gcc/fortran/dump-parse-tree.cc b/gcc/fortran/dump-parse-tree.cc index 68122e3e6fd..0e888fafe7b 100644 --- a/gcc/fortran/dump-parse-tree.cc +++ b/gcc/fortran/dump-parse-tree.cc @@ -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; diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index 9a00e6dea6f..a8667a6e6d3 100644 --- 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) op:8; + bool readonly; + } map; gfc_expr *align; struct { diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index 2952cd300ac..af769d9efbd 100644 --- 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; } @@ -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 " diff --git a/gcc/fortran/trans-decl.cc b/gcc/fortran/trans-decl.cc index b0fd25e92a3..1ff1dda026a 100644 --- a/gcc/fortran/trans-decl.cc +++ b/gcc/fortran/trans-decl.cc @@ -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: diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index cf741cebf91..a4628e460bd 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -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; 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..171f96c08db --- /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" } } */ + + + 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..069fec0a0d5 --- /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" } } + + + diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc index 25d191b10fd..9604c3eecc5 100644 --- 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; diff --git a/gcc/tree.h b/gcc/tree.h index 4c04245e2b1..1301491587f 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -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) \