From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id E32283858D20; Thu, 26 Oct 2023 09:43:21 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org E32283858D20 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com ARC-Filter: OpenARC Filter v1.0.0 sourceware.org E32283858D20 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=68.232.141.98 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1698313404; cv=none; b=JkB5Z7UV7GnWN4VnDXrO1lidTbGfyjLEvKQrHMs1rR2BsKE/9rA7vls6ud23CmbylRTxGvkkRzK3zxRYdadyTdnvbsF3VN3+l7HqekY19tRu5muaiRa3u46KlsRPON2VE3/h8FPVq0h0TOzw4qLCfSfFryKiNUzROPNZourkuTE= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1698313404; c=relaxed/simple; bh=c2V3oMlLDyVqa59Ft6MB02MwyKQGKEBUdoqysl30UV8=; h=From:To:Subject:Date:Message-ID:MIME-Version; b=X5eCsF5bBzN7ZgIkI9l0Auaq0fm7CbrDdrTFziAXX/E30RFrZptYXASiYuAi6I09RwdkmoNvCRye73SSdIT6/GJQ73GGQrZ9DTdmwDLe0Gzd4YSIGsemF3ObzNGIyym1nl/H8I2QR0NfBJVfHDJ2C7+pJkWuphDjiwewf0thRxU= ARC-Authentication-Results: i=1; server2.sourceware.org X-CSE-ConnectionGUID: bUGnOVvzRy+iETM8wwmVkA== X-CSE-MsgGUID: cpu7tlswTFmGRN8mPaDaVw== X-IronPort-AV: E=Sophos;i="6.03,253,1694764800"; d="scan'208";a="23353793" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa2.mentor.iphmx.com with ESMTP; 26 Oct 2023 01:43:20 -0800 IronPort-SDR: f0KAi4BXVSNQMQ2ogFRm1fu6IKyL6HfmgJaNrpvhk5xzzETSMK2GWSlrfVN3kFAiR+yedZv38I fFL6GFRpSL6Ad2QTgU+a/MfrMCi3cjZZL3OrKZI7DfYYTomAK38ecP3X6JspcxBMs84oUDIFL4 bICiohhxmXaLyuX5giEtN2vDk1SMslr9F2BnAmquXhBmYeWmxZE/6a9bqIQXEjHmbPEIHWTHVD aZy3zwev8JtiCkF+rmH0tVWZX57xH4Jm6+hQVfyUpD7ZKEwStSwGKKwgdbGfR2y++Ryopm6W1u lx0= From: Thomas Schwinge To: Chung-Lin Tang , Tobias Burnus CC: , Catherine Moore , Subject: Re: [PATCH, OpenACC 2.7, v2] readonly modifier support in front-ends In-Reply-To: References: <87lefaaesb.fsf@euler.schwinge.homeip.net> User-Agent: Notmuch/0.29.3+94~g74c3f1b (https://notmuchmail.org) Emacs/28.2 (x86_64-pc-linux-gnu) Date: Thu, 26 Oct 2023 11:43:13 +0200 Message-ID: <87ttqd7m8e.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 Content-Type: text/plain; charset="utf-8" Content-Transfer-Encoding: quoted-printable X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-13.mgc.mentorg.com (139.181.222.13) To svr-ies-mbx-10.mgc.mentorg.com (139.181.222.10) X-Spam-Status: No, score=-5.8 required=5.0 tests=BAYES_00,HEADER_FROM_DIFFERENT_DOMAINS,KAM_DMARC_STATUS,SPF_HELO_PASS,SPF_PASS,TXREP autolearn=no autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org List-Id: Hi! On 2023-08-07T21:58:27+0800, Chung-Lin Tang wro= te: > 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 =3D false) >>>> + tree list, bool allow_deref =3D false, >>>> + bool *readonly =3D 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 plac= es >>> 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 direc= tly > 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'. = =E2=80=93 >> this seems to be cleaner. > > I've adjusted 'u.map' to be a named struct now, and updated the reference= s. 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 :") =3D=3D 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 th= at >>> 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 nelsas= y 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, o= r >>> 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' i= s >>> 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 s= ense 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 =3D=3D 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 TR= EE_READONLY? | | int a; | #pragma acc parallel copyin(readonly : a) | {...} | a =3D 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 =E2=80=93 but it is still ugly and, fo= r | 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 =3D 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_M= AP_READONLY > to mark the clauses of "base-pointers of array-sections" as OMP_CLAUSE_MA= P_POINTS_TO_READONLY, > and later this gradually gets relayed to alias oracle routines in tree-ss= a-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, pr= agma_omp_clause c_kind, > default: > gcc_unreachable (); > } > - tree nl, c; > - nl =3D c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list, tru= e); > > - for (c =3D nl; c !=3D list; c =3D OMP_CLAUSE_CHAIN (c)) > - OMP_CLAUSE_SET_MAP_KIND (c, kind); > + tree nl =3D list; > + bool readonly =3D false; > + matching_parens parens; > + if (parens.require_open (parser)) > + { > + /* Turn on readonly modifier parsing for copyin clause. */ > + if (c_kind =3D=3D PRAGMA_OACC_CLAUSE_COPYIN) > + { > + c_token *token =3D c_parser_peek_token (parser); > + if (token->type =3D=3D CPP_NAME > + && !strcmp (IDENTIFIER_POINTER (token->value), "readonly") > + && c_parser_peek_2nd_token (parser)->type =3D=3D CPP_COLON) > + { > + c_parser_consume_token (parser); > + c_parser_consume_token (parser); > + readonly =3D true; > + } > + } > + location_t loc =3D 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 =3D c_parser_omp_variable_list (parser, loc, OMP_CLAUSE_MAP, li= st, true); > + parens.skip_until_found_close (parser); > + } > + > + for (tree c =3D nl; c !=3D list; c =3D OMP_CLAUSE_CHAIN (c)) > + { > + OMP_CLAUSE_SET_MAP_KIND (c, kind); > + if (readonly) > + OMP_CLAUSE_MAP_READONLY (c) =3D 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 =3D NULL_TREE; > + bool readonly =3D false; > + matching_parens parens; > + > + if (parens.require_open (parser)) > + { > + c_token *token =3D c_parser_peek_token (parser); > + if (token->type =3D=3D CPP_NAME > + && !strcmp (IDENTIFIER_POINTER (token->value), "readonly") > + && c_parser_peek_2nd_token (parser)->type =3D=3D CPP_COLON) > + { > + c_parser_consume_token (parser); > + c_parser_consume_token (parser); > + readonly =3D true; > + } > + location_t loc =3D c_parser_peek_token (parser)->location; Similar. (That is, here, location of the directive.) > + clauses =3D c_parser_omp_variable_list (parser, loc, OMP_CLAUSE__C= ACHE_, > + NULL_TREE); > + parens.skip_until_found_close (parser); > + } > + > + if (readonly) > + for (tree c =3D clauses; c; c =3D OMP_CLAUSE_CHAIN (c)) > + OMP_CLAUSE__CACHE__READONLY (c) =3D 1; > > - clauses =3D c_parser_omp_var_list_parens (parser, OMP_CLAUSE__CACHE_, = NULL); > clauses =3D 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) : om= p_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 reado= nly =3D false) > { > gfc_omp_namelist **head =3D 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 =3D *head; n; n =3D n->next) > - n->u.map_op =3D map_op; > + { > + n->u.map.op =3D map_op; > + n->u.map.readonly =3D 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=C3=BC=C3=9Fe Thomas ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstra=C3=9Fe 201= , 80634 M=C3=BCnchen; Gesellschaft mit beschr=C3=A4nkter Haftung; Gesch=C3= =A4ftsf=C3=BChrer: Thomas Heurung, Frank Th=C3=BCrauf; Sitz der Gesellschaf= t: M=C3=BCnchen; Registergericht M=C3=BCnchen, HRB 106955