From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa1.mentor.iphmx.com (esa1.mentor.iphmx.com [68.232.129.153]) by sourceware.org (Postfix) with ESMTPS id 5B1023858CDB for ; Thu, 20 Jul 2023 15:08:38 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 5B1023858CDB Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="6.01,219,1684828800"; d="scan'208";a="13936460" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa1.mentor.iphmx.com with ESMTP; 20 Jul 2023 07:08:36 -0800 IronPort-SDR: s/n7mmpmDOxNrJtM33UjkZ1vZ5LgKebLa1qENno0McePIG0YxRyb5OZAD9cfCcCcDISEQZ7//0 kctkpM5iMTNFus/BSCODrR48U5VEZ3EmJrf2k/a8Urrf6mYq/HhSu1+UlAxHLXQjzs7Ehw9j2z k8Z5TuNFxm29WMqSr130PpacKV0thnLg3Ssq4HL5gqww0a7WmIAIVPfvDErbP3t3hcFdt77nV+ 98P4JAiSr0J/6tUbQwGkPGjUtqNeCtFnfeDz1QtzSCdaB+eXj5wq8GEbulWlOcrkLC79l6Bldy siY= Message-ID: Date: Thu, 20 Jul 2023 17:08:31 +0200 MIME-Version: 1.0 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:102.0) Gecko/20100101 Thunderbird/102.13.0 Subject: Re: [PATCH, OpenACC 2.7] readonly modifier support in front-ends Content-Language: en-US To: Thomas Schwinge , Chung-Lin Tang CC: , Catherine Moore References: <87lefaaesb.fsf@euler.schwinge.homeip.net> From: Tobias Burnus In-Reply-To: <87lefaaesb.fsf@euler.schwinge.homeip.net> Content-Type: text/plain; charset="UTF-8"; format=flowed Content-Transfer-Encoding: quoted-printable X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-11.mgc.mentorg.com (139.181.222.11) To svr-ies-mbx-12.mgc.mentorg.com (139.181.222.12) X-Spam-Status: No, score=-5.4 required=5.0 tests=BAYES_00,HEADER_FROM_DIFFERENT_DOMAINS,KAM_DMARC_STATUS,NICE_REPLY_A,SPF_HELO_PASS,SPF_PASS,TXREP,T_SCC_BODY_TEXT_LINE 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 Thomas & Chung-Lin, On 20.07.23 15:33, Thomas Schwinge wrote: > On 2023-07-11T02:33:58+0800, Chung-Lin Tang > wrote: >> +++ b/gcc/c/c-parser.cc >> @@ -14059,7 +14059,8 @@ c_parser_omp_variable_list (c_parser *parser, >> >> static tree >> c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code k= ind, >> - 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 places > where it's used. I concur. The same issue also occurred for OpenMP's c_parser_omp_clause_to, and c_parser_omp_clause_from and the 'present' modifier. For it, I created a combined function but the main reason for that is that OpenMP also permits more modifiers (like 'iterators'), which would cause more duplication of code ('iterator' is not yet supported). For something as simple to parse as this modifier, I would just do it at the two places =E2=80=93 as Thomas suggested. >> +++ b/gcc/fortran/gfortran.h >> @@ -1360,7 +1360,11 @@ typedef struct gfc_omp_namelist >> { >> gfc_omp_reduction_op reduction_op; >> gfc_omp_depend_doacross_op depend_doacross_op; >> - gfc_omp_map_op map_op; >> + struct >> + { >> + ENUM_BITFIELD (gfc_omp_map_op) map_op:8; >> + bool readonly; >> + }; >> gfc_expr *align; >> struct >> { > [...] Thus, the above looks good to me. I concur but I wonder whether it would be cleaner to name the struct; this makes it also more obvious what belongs together in the union. Namely, naming the struct 'map' and then changing the 45 users from 'u.map_op' to 'u.map.op' and the new 'u.readonly' to 'u.map.readonly'. =E2= =80=93 this seems to be cleaner. >> static bool >> gfc_match_omp_map_clause (gfc_omp_namelist **list, gfc_omp_map_op map_= op, >> - bool allow_common, bool allow_derived) >> + bool allow_common, bool allow_derived, bool read= only =3D false) >> { > Similar to 'c_parser_omp_var_list_parens' above, I concur that not doing it here is cleaner. > again, for > example (random), like 'ancestor :', or 'conditional :' are parsed -- > which you're mostly already doing I think OpenMP's "present" (as modifier to "omp target updates"'s "to"/"from") is a better example than "ancestor" as for present we also have a list. See: gfc_match_motion_var_list how to handle the headp. (There an extra functions was used as in the future also other modifiers like 'iterator' will be used.) However, as Thomas noted, the patch contains also an example (see further down in Thomas' email, not quoted here). > Or, we could add a new 'gcc/fortran/gfortran.h:gfc_omp_map_op' item > 'OMP_MAP_TO_READONLY', which eventually translates into 'OMP_MAP_TO' with > 'readonly' set? I think having the additional flag is easier to understand - and at least memory wise we do not save memory as it is in a union. The advantage of not having a union is that accessing the int-enum is faster than accessi= ng an char-wide bitset enum. In terms of code changes (and without having a closer look), the two approaches seems to be be similar. Hence, using OMP_MAP_TO_READONLY for OpenACC would be fine, too. And I do not have a strong preference for either. * * * I did wonder about the following, but I now believe it won't affect the choice. Namely, we want to handle at some point the following: !$omp target firstprivate(var) allocator(omp_const_mem_alloc: var) This could be turned into GOMP_MAP_FIRSTPRIVATE... + OMP_.*READONLY flag. But if we don't do it in the FE, the internal Fortran representation does not matter. Advantage for doing it in the ME: Only one code location, especially as we might use the opportunity to also check that the omp_const_mem_alloc is only used with privatization (in OpenMP). Difference: OpenMP uses 'firstprivate' (i.e. private copy, no reference cou= nt bump, only permitted for 'target') while OpenACC uses 'copy' which implies refere= nce counting and permitted in 'acc (enter/exit) data' and not only for compute = constructs. OpenMP in principle also permits user-defined allocator with a constant memory space - I am not completely sure whether/when it can be used with omp target firstprivate(...) allocator(my_alloc : ...) > Then we'd just here call the (unaltered) > 'gfc_match_omp_map_clause', with > 'readonly ? OMP_MAP_TO_READONLY : OMP_MAP_TO'? Per > 'git grep --cached '[^G]OMP_MAP_TO[^F]' -- gcc/fortran/' not a lot of > places need adjusting for that (most of the 'gcc/fortran/openmp.cc' ones > are not applicable). I think either would work. =E2=80=93 I have no strong feeling what's better= . But you still need to handle it for clause resolution. > + 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 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 per= mit: !$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 n= c co p y(a)".) A " " matches zero or more whitespaces, but with gfc_match_space you can fi= nd out whether there was whitespace or not. Whether the tailing " " in the gfc_match matters or not, depends on what co= mes next. If there is a "gfc_gobble_whitespace ();", everything is fine. If not, the = next to match has to start with a " ", which is usually ugly; an exception is " , " or " = )" which still is somewhat fine. I think that it is mostly implemented correctly, but I wouldn't be surprise= d if a space is missing in some matches - be it a tailing white space or e.g. in "= foo:" before the colon. BTW: One reason of stripping tailing spaces before matching a non-whitespac= e: the associated location is the one before the parsing; thus, for a match error = or when saving the old_locus, pointing to the first non-whitespace looks nicer than pointi= ng to the (first of the) whitspace character(s). >> + if (readonly) >> + for (gfc_omp_namelist *n =3D *head; n; n =3D n->next) >> + n->u.readonly =3D true; > This already looks like how I thought it should look like. Indeed.--- a/gcc/tree.h >> +++ b/gcc/tree.h >> @@ -1813,6 +1813,14 @@ class auto_suppress_location_wrappers >> #define OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE(NODE) \ >> (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.addressable_f= lag) >> >> +/* Nonzero if OpenACC 'readonly' modifier set, used for 'copyin'. */ >> +#define OMP_CLAUSE_MAP_READONLY(NODE) \ >> + TREE_READONLY (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)) >> + >> +/* Same as above, for use in OpenACC cache directives. */ >> +#define OMP_CLAUSE__CACHE__READONLY(NODE) \ >> + TREE_READONLY (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE__CACHE_)) > I'm not sure if these special accessor functions are actually useful, or > we should just directly use 'TREE_READONLY' instead? We're only using > them in contexts where it's clear that the 'OMP_CLAUSE_SUBCODE_CHECK' is > satisfied, for example. I find directly using TREE_READONLY confusing. > Setting 'TREE_READONLY' of the 'OMP_CLAUSE_DECL' instead of the clause > itself isn't the right thing to do -- or is it, and might already > indicate to the middle end the desired semantics? But does it maybe > conflict with front end/language-level use of 'TREE_READONLY' for 'const' > etc. (I suppose), and thus diagnostics for mismatches? I think is is cleaner not to one flag to mean two different things. In particular, wouldn't the following cause issues, if you mark 'a' as TREE= _READONLY? int a; #pragma acc parallel copyin(readonly : a) {...} a =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, for declarations, unshare_expr does not create a copy! Tobias ----------------- 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