Hi Jakub, On 2021/6/24 11:55 PM, Jakub Jelinek wrote: > On Fri, May 14, 2021 at 09:20:25PM +0800, Chung-Lin Tang wrote: >> diff --git a/gcc/gimplify.c b/gcc/gimplify.c >> index e790f08b23f..69c4a8e0a0a 100644 >> --- a/gcc/gimplify.c >> +++ b/gcc/gimplify.c >> @@ -10374,6 +10374,7 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) >> gcc_unreachable (); >> } >> OMP_CLAUSE_SET_MAP_KIND (clause, kind); >> + OMP_CLAUSE_MAP_IMPLICIT_P (clause) = 1; >> if (DECL_SIZE (decl) >> && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST) >> { > > As Thomas mentioned, there is now also OMP_CLAUSE_MAP_IMPLICIT that means > something different: > /* Nonzero on map clauses added implicitly for reduction clauses on combined > or composite constructs. They shall be removed if there is an explicit > map clause. */ > Having OMP_CLAUSE_MAP_IMPLICIT and OMP_CLAUSE_MAP_IMPLICIT_P would be too > confusing. So either we need to use just one flag for both purposes or > have two different flags and find a better name for one of them. > The former would be possible if no OMP_CLAUSE_MAP clauses added by the FEs > are implicit - then you could clear OMP_CLAUSE_MAP_IMPLICIT in > gimplify_scan_omp_clauses. I wonder if it is the case though, e.g. doesn't > your "Improve OpenMP target support for C++ [PR92120 v4]" patch add a lot of > such implicit map clauses (e.g. the this[:1] and various others)? I have changed the name to OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P, to signal that this bit is to be passed to the runtime. Right now its intended to be used by clauses created by the middle-end, but front-end uses like that for C++ could be clarified later. > Also, gimplify_adjust_omp_clauses_1 sometimes doesn't add just one map > clause, but several, shouldn't those be marked implicit too? And similarly > it calls lang_hooks.decls.omp_finish_clause which can add even further map > clauses implicitly, shouldn't those be implicit too (in that case copy > the flag from the clause it is called on to the extra clauses it adds)? > > Also as Thomas mentioned, it should be restricted to non-OpenACC, > it can check gimplify_omp_ctxp->region_type if it is OpenMP or OpenACC. Agreed, I've adjusted the patch to only to this implicit setting for OpenMP. This reduces a lot of the originally needed scan test adjustment for existing OpenACC testcases. >> @@ -10971,9 +10972,15 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, >> list_p = &OMP_CLAUSE_CHAIN (c); >> } >> >> - /* Add in any implicit data sharing. */ >> + /* Add in any implicit data sharing. Implicit clauses are added at the start > > Two spaces after dot in comments. Done. >> + of the clause list, but after any non-map clauses. */ >> struct gimplify_adjust_omp_clauses_data data; >> - data.list_p = list_p; >> + tree *implicit_add_list_p = orig_list_p; >> + while (*implicit_add_list_p >> + && OMP_CLAUSE_CODE (*implicit_add_list_p) != OMP_CLAUSE_MAP) >> + implicit_add_list_p = &OMP_CLAUSE_CHAIN (*implicit_add_list_p); > > Why are the implicit map clauses added first and not last? As I also explained in the first submission email, due to the processing order, if implicit classes are added last (and processed last), for example: #pragma omp target map(tofrom: var.ptr[:N]) map(tofrom: var[implicit]) { // access of var.ptr[] } The explicit var.ptr[:N] will not find anything to map, because the (implicit) map(var) has not been seen yet, and the assumed array section attachment behavior will fail. Only an order like: map(tofrom: var[implicit]) map(tofrom: var.ptr[:N]) will the usual assumed behavior show. And yes, this depends on the new behavior implemented by patch [1], which I still need you to review. e.g. for map(var.ptr[:N]), the proper behavior should *only* map the array section but NOT the base-pointer. [1] https://gcc.gnu.org/pipermail/gcc-patches/2021-May/571195.html > There is also the OpenMP 5.1 [352:17-22] case which basically says that the > implicit mappings should be ignored if there are explicit ones on the same > construct (though, do we really create implicit clauses in that case?). Implicit clauses do not appear to be created if there's an explicit clause already existing. >> +#define GOMP_MAP_IMPLICIT (GOMP_MAP_FLAG_SPECIAL_3 \ >> + | GOMP_MAP_FLAG_SPECIAL_4) >> +/* Mask for entire set of special map kind bits. */ >> +#define GOMP_MAP_FLAG_SPECIAL_BITS (GOMP_MAP_FLAG_SPECIAL_0 \ >> + | GOMP_MAP_FLAG_SPECIAL_1 \ >> + | GOMP_MAP_FLAG_SPECIAL_2 \ >> + | GOMP_MAP_FLAG_SPECIAL_3 \ >> + | GOMP_MAP_FLAG_SPECIAL_4) ... >> +#define GOMP_MAP_IMPLICIT_P(X) \ >> + (((X) & GOMP_MAP_FLAG_SPECIAL_BITS) == GOMP_MAP_IMPLICIT) > > I think here we need to decide with which GOMP_MAP* kinds the implicit > bit will need to be combined with, with looking forward into what features > we still need to implement for OpenMP 5.0/5.1 (not aware of anything in 5.2 > that would need special care but perhaps I've missed it). > > E.g. for declare mapper those single OMP_CLAUSE_MAPs with the implicit > bit might need to be split into various smaller ones, map this FIELD_DECL, > map that other FIELD_DECL, what it points to, etc. Even without declare > mapper, the spec now says that mapping a structure is as if each member is > mapped separately and 5.2 is going to say there is a predefined declare > mapper that does that (which is of course something we don't want under the > hood, we don't want thousands of maps, but virtually split it on a field by > field basis, do all the special stuff - e.g. how pointers are to be mapped > as zero array sections with pointer attachment, how C++ references are to be > handled etc. and then virtually coalesce all the adjacent fields that have > the same treatment back). I think that actually means we should defer most > of the map struct etc. handling we do in gimplify_scan_omp_clauses, instead > note the explicit map clauses and what they refer to, add implicit ones (and > set implicit bit on those), then go over all clauses (explicit and > implicit), do the declare mapper processing, do the sorting of what that > produces based on base expressions/pointers etc. > At this point, I'm not sure if GOMP_MAP_IMPLICIT can or can't appear > together e.g. with GOMP_MAP_STRUCT. Currently with the place of setting OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P in gimplify_adjust_omp_clauses_1, only things like GOMP_MAP_TO/FROM/FORCE_PRESENT/etc. be set. I already had some trouble designating how GOMP_MAP_IMPLICIT would be encoded within the current bits. My guess is that with more sophisticated features like declare mappers, the whole interface probably needs to be extended somewhat. >> @@ -405,8 +422,24 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep, >> static int >> get_kind (bool short_mapkind, void *kinds, int idx) >> { >> - return short_mapkind ? ((unsigned short *) kinds)[idx] >> - : ((unsigned char *) kinds)[idx]; >> + int val = (short_mapkind >> + ? ((unsigned short *) kinds)[idx] >> + : ((unsigned char *) kinds)[idx]); >> + >> + if (GOMP_MAP_IMPLICIT_P (val)) >> + val &= ~GOMP_MAP_IMPLICIT; > > As the particular bit isn't used for anything right now, > perhaps just do val &= ~GOMP_MAP_IMPLICIT unconditionally? GOMP_MAP_IMPLICIT is (GOMP_MAP_FLAG_SPECIAL_3 | GOMP_MAP_FLAG_SPECIAL_4), and those two bits are also partially used for some other things. So an specific check should be needed. > But, on the other side, for !short_mapkind you do not want > to mask that bit out, only the low 3 bits are the mapping > type and the upper bits are log2 of alignment, so the > above for !short_mapkind would in weird way change some > alignments. I see, I've added a short_mapkind check in get_kind. >> + return val; >> +} >> + >> + >> +static bool >> +get_implicit (bool short_mapkind, void *kinds, int idx) >> +{ >> + int val = (short_mapkind >> + ? ((unsigned short *) kinds)[idx] >> + : ((unsigned char *) kinds)[idx]); >> + >> + return GOMP_MAP_IMPLICIT_P (val); > > Similarly can return 0 for !short_mapkind, the compatibility GOMP_target etc. > APIs will never have those in. This too added. Re-tested for trunk without regressions, is this okay now? Thanks, Chung-Lin 2021-11-05 Chung-Lin Tang include/ChangeLog: * gomp-constants.h (GOMP_MAP_FLAG_SPECIAL_3): Define special bit macro. (GOMP_MAP_IMPLICIT): New special map kind bits value. (GOMP_MAP_FLAG_SPECIAL_BITS): Define helper mask for whole set of special map kind bits. (GOMP_MAP_IMPLICIT_P): New predicate macro for implicit map kinds. gcc/ChangeLog: * tree.h (OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P): New access macro for 'implicit' bit, using 'base.deprecated_flag' field of tree_node. * tree-pretty-print.c (dump_omp_clause): Add support for printing implicit attribute in tree dumping. * gimplify.c (gimplify_adjust_omp_clauses_1): Set OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P to 1 if map clause is implicitly created. (gimplify_adjust_omp_clauses): Adjust place of adding implicitly created clauses, from simple append, to starting of list, after non-map clauses. * omp-low.c (lower_omp_target): Add GOMP_MAP_IMPLICIT bits into kind values passed to libgomp for implicit maps. gcc/testsuite/ChangeLog: * c-c++-common/gomp/target-implicit-map-1.c: New test. * c-c++-common/goacc/combined-reduction.c: Adjust scan test pattern. * c-c++-common/goacc/firstprivate-mappings-1.c: Likewise. * c-c++-common/goacc/mdc-1.c: Likewise. * g++.dg/goacc/firstprivate-mappings-1.C: Likewise. libgomp/ChangeLog: * target.c (gomp_map_vars_existing): Add 'bool implicit' parameter, add implicit map handling to allow a "superset" existing map as valid case. (get_kind): Adjust to filter out GOMP_MAP_IMPLICIT bits in return value. (get_implicit): New function to extract implicit status. (gomp_map_fields_existing): Adjust arguments in calls to gomp_map_vars_existing, and add uses of get_implicit. (gomp_map_vars_internal): Likewise. * testsuite/libgomp.c-c++-common/target-implicit-map-1.c: New test.