On Wed, 3 Apr 2024, Chung-Lin Tang wrote:
> Hi Richard, Thomas,
>
> On 2023/10/30 8:46 PM, Richard Biener wrote:
> >>
> >> What Chung-Lin's first patch does is mark the OMP clause for 'x' (not the
> >> 'x' decl itself!) as 'readonly', via a new 'OMP_CLAUSE_MAP_READONLY'
> >> flag.
> >>
> >> The actual optimization then is done in this second patch. Chung-Lin
> >> found that he could use 'SSA_NAME_POINTS_TO_READONLY_MEMORY' for that.
> >> I don't have much experience with most of the following generic code, so
> >> would appreciate a helping hand, whether that conceptually makes sense as
> >> well as from the implementation point of view:
>
> First of all, I have removed all of the gimplify-stage scanning and setting of
> DECL_POINTS_TO_READONLY and SSA_NAME_POINTS_TO_READONLY_MEMORY (so no changes
> to
> gimplify.cc now)
>
> I remember this code was an artifact of earlier attempts to allow
> struct-member
> pointer mappings to also work (e.g. map(readonly:rec.ptr[:N])), but failed
> anyways.
> I think the omp_data_* member accesses when building child function side
> receiver_refs is blocking points-to analysis from working (didn't try digging
> deeper)
>
> Also during gimplify, VAR_DECLs appeared to be reused (at least in some
> cases) for map
> clause decl reference building, so hoping that the variables "happen to be"
> single-use and
> DECL_POINTS_TO_READONLY relaying into SSA_NAME_POINTS_TO_READONLY_MEMORY does
> appear to be
> a little risky.
>
> However, for firstprivate pointers processed during omp-low, it appears to be
> somewhat different.
> (see below description)
>
> > No, I don't think you can use that flag on non-default-defs, nor
> > preserve it on copying. So
> > it also doesn't nicely extend to DECLs as done by the patch. We
> > currently _only_ use it
> > for incoming parameters. When used on arbitrary code you can get to for
> > example
> >
> > ptr1(points-to-readony-memory) = &p->x;
> > ... access via ptr1 ...
> > ptr2 = &p->x;
> > ... access via ptr2 ...
> >
> > where both are your OMP regions differently constrained (the constrain is
> > on the
> > code in the region, _not_ on the actual protections of the pointed to
> > data, much like
> > for the fortran case). But now CSE comes along and happily replaces all
> > ptr2
> > with ptr2 in the second region and ... oops!
>
> Richard, I assume what you meant was "happily replaces all ptr2 with ptr1 in
> the second region"?
>
> That doesn't happen, because during omp-lower/expand, OMP target regions
> (which is all that
> this applies currently) is separated into different individual child
> functions.
>
> (Currently, the only "effective" use of DECL_POINTS_TO_READONLY is during
> omp-lower, when
> for firstprivate pointers (i.e. 'a' here) we set this bit when constructing
> the first load
> of this pointer)
>
> #pragma acc parallel copyin(readonly: a[:32]) copyout(r)
> {
> foo (a, a[8]);
> r = a[8];
> }
> #pragma acc parallel copyin(readonly: a[:32]) copyout(r)
> {
> foo (a, a[12]);
> r = a[12];
> }
>
> After omp-expand (before SSA):
>
> __attribute__((oacc parallel, omp target entrypoint, noclone))
> void main._omp_fn.1 (const struct .omp_data_t.3 & restrict .omp_data_i)
> {
> ...
> <bb 5> :
> D.2962 = .omp_data_i->D.2947;
> a.8 = D.2962;
So 'readonly: a[:32]' is put in .omp_data_i->D.2947 in the caller
and extracted here. And you arrange for 'a.8' to have
DECL_POINTS_TO_READONLY set by "magic"? Looking at this I wonder
if it would be more useful to "const qualify" (but "really", not
in the C sense) .omp_data_i->D.2947 instead? Thus have a
FIELD_POINTS_TO_READONLY_MEMORY flag on the FIELD_DECL.
Points-to analysis should then be able to handle this similar to how
it handles loads of restrict qualified pointers. Well, of course not
as simple since it now adds "qualifiers" to storage since I presume
the same object can be both readonly and not readonly like via
#pragma acc parallel copyin(readonly: a[:32], a[33:64]) copyout(r)
? That is, currently there's only one "readonly" object kind in
points-to, that's STRING_CSTs which get all globbed to string_id
and "ignored" for alias purposes since you can't change them.
So possibly you want to combine this with restrict qualifying the
pointer so we know there's no other (read-write) access to the memory
possible. But then you might get all the good stuff already by
_just_ doing that restrict qualification and ignoring the readonly-ness?
> r.1 = (*a.8)[12];
> foo (a.8, r.1);
> r.1 = (*a.8)[12];
> D.2965 = .omp_data_i->r;
> *D.2965 = r.1;
> return;
> }
>
> __attribute__((oacc parallel, omp target entrypoint, noclone))
> void main._omp_fn.0 (const struct .omp_data_t.2 & restrict .omp_data_i)
> {
> ...
> <bb 3> :
> D.2968 = .omp_data_i->D.2939;
> a.4 = D.2968;
> r.0 = (*a.4)[8];
> foo (a.4, r.0);
> r.0 = (*a.4)[8];
> D.2971 = .omp_data_i->r;
> *D.2971 = r.0;
> return;
> }
>
> So actually, the creating of DECL_POINTS_TO_READONLY and its relaying to
> SSA_NAME_POINTS_TO_READONLY_MEMORY here, is actually quite similar to a
> default-def
> for an PARM_DECL, at least conceptually.
>
> (If offloading was structured significantly differently, say if child
> functions
> were separated much earlier before omp-lowering, than this readonly-modifier
> might
> possibly be a direct application of 'r' in the "fn spec" attribute)
>
> Other changes since first version of patch include:
> 1) update of C/C++ FE changes to new style in c-family/c-omp.cc
> 2) merging of two if cases in fortran/trans-openmp.cc like Thomas suggested
> 3) Update of readonly-2.c testcase to scan before/after "fre1" pass, to
> verify removal of a MEM load, also as Thomas suggested.
>
> I have re-tested this patch using mainline, with no regressions. Is this
> okay for mainline?
+/* In a VAR_DECL, set for variables regarded as pointing to memory not
written
+ to. SSA_NAME_POINTS_TO_READONLY_MEMORY gets set for SSA_NAMEs created
from
+ such VAR_DECLs. Currently used by OpenACC 'readonly' modifier in
copyin
+ clauses. */
+#define DECL_POINTS_TO_READONLY(NODE) \
+ (TREE_CHECK (NODE, VAR_DECL)->decl_common.decl_not_flexarray)
you need to document uses of flags in tree-core.h to avoid clashes.
Also since this doesn't apply to all DECLs it should be named
VAR_POINTS_TO_...
I still think this is too fragile - there's no real constraints
on what VAR_DECL we create SSA names off, so the automatism
in make_ssa_name_fn and esp. copy_var_decl and via copy_node
copy_decl_no_change, thus during inlining, makes your arguments
only apply to the use for OpenMP - but nothing above hints at
this is just usable there, asking for trouble.
Sorry for the delay,
Richard.
> Thanks,
> Chung-Lin
>
> 2024-04-03 Chung-Lin Tang <clt...@baylibre.com>
>
> gcc/c-family/ChangeLog:
>
> * c-omp.cc (c_omp_address_inspector::expand_array_base):
> Set OMP_CLAUSE_MAP_POINTS_TO_READONLY on pointer clause.
> (c_omp_address_inspector::expand_component_selector): Likewise.
>
> gcc/fortran/ChangeLog:
>
> * trans-openmp.cc (gfc_trans_omp_array_section):
> Set OMP_CLAUSE_MAP_POINTS_TO_READONLY on pointer clause.
>
> gcc/ChangeLog:
>
> * gimple-expr.cc (copy_var_decl): Copy DECL_POINTS_TO_READONLY
> for VAR_DECLs.
> * omp-low.cc (lower_omp_target): Set DECL_POINTS_TO_READONLY for
> variables of receiver refs.
> * tree-pretty-print.cc (dump_omp_clause):
> Print OMP_CLAUSE_MAP_POINTS_TO_READONLY.
> (dump_generic_node): Print SSA_NAME_POINTS_TO_READONLY_MEMORY.
> * tree-ssanames.cc (make_ssa_name_fn): Set
> SSA_NAME_POINTS_TO_READONLY_MEMORY if DECL_POINTS_TO_READONLY is set.
> * tree.h (DECL_POINTS_TO_READONLY): New macro.
> (OMP_CLAUSE_MAP_POINTS_TO_READONLY): New macro.
>
> gcc/testsuite/ChangeLog:
>
> * c-c++-common/goacc/readonly-1.c: Adjust testcase.
> * c-c++-common/goacc/readonly-2.c: New testcase.
> * gfortran.dg/goacc/readonly-1.f90: Adjust testcase.
>
>
>
>
>
>
>
>
>
>
--
Richard Biener <rguent...@suse.de>
SUSE Software Solutions Germany GmbH,
Frankenstrasse 146, 90461 Nuernberg, Germany;
GF: Ivo Totev, Andrew McDonald, Werner Knoblich; (HRB 36809, AG Nuernberg)
