On Fri, Oct 27, 2023 at 4:28 PM Thomas Schwinge <tho...@codesourcery.com> wrote:
>
> Hi!
>
> Richard, as the original author of 'SSA_NAME_POINTS_TO_READONLY_MEMORY':
> 2018 commit 6214d5c7e7470bdd5ecbeae668c2522551bfebbc (Subversion r263958)
> "Move const_parm trick to generic code"; 'gcc/tree.h':
>
> /* Nonzero if this SSA_NAME is known to point to memory that may not
> be written to. This is set for default defs of function parameters
> that have a corresponding r or R specification in the functions
> fn spec attribute. This is used by alias analysis. */
> #define SSA_NAME_POINTS_TO_READONLY_MEMORY(NODE) \
> SSA_NAME_CHECK (NODE)->base.deprecated_flag
>
> ..., may I ask you to please help review the following patch
> (full-quoted)?
>
> For context: this patch here ("second patch") depends on a first patch:
> <inbox.sourceware.org/d0e6013f-ca38-b98d-dc01-b30adbd59...@siemens.com>
> "[PATCH, OpenACC 2.7] readonly modifier support in front-ends". That one
> is still under review/rework; so you're not able to apply this second
> patch here.
>
> In a nutshell: a 'readonly' modifier has been added to the OpenACC
> 'copyin' clause (copy host to device memory, don't copy back at end of
> region):
>
> | If the optional 'readonly' modifier appears, then the implementation may
> assume that the data
> | referenced by _var-list_ is never written to within the applicable region.
>
> That is, for example (untested):
>
> #pragma acc routine
> void escape(int *);
>
> int x[32] = [...];
> #pragma acc parallel copyin(readonly: x)
> {
> int a1 = x[3];
> escape(x);
> int a2 = x[3]; // Per 'readonly', don't need to reload 'x[3]' here.
> //x[22] = 0; // Invalid -- but no diagnostic mandated.
> }
>
> 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:
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!
So no, re-using SSA_NAME_POINTS_TO_READONLY_MEMORY doesn't look good.
Richard.
> On 2023-07-25T23:52:06+0800, Chung-Lin Tang via Gcc-patches
> <gcc-patches@gcc.gnu.org> wrote:
> > On 2023/7/11 2:33 AM, Chung-Lin Tang via Gcc-patches wrote:
> >> As we discussed earlier, the work for actually linking this to middle-end
> >> points-to analysis is a somewhat non-trivial issue. This first patch allows
> >> the language feature to be used in OpenACC directives first (with no
> >> effect for now).
> >> The middle-end changes are probably going to be a later patch.
> >
> > This second patch tries to link the readonly modifier to points-to analysis.
> >
> > There already exists SSA_NAME_POINTS_TO_READONLY_MEMORY and it's support in
> > the
> > alias oracle routines in tree-ssa-alias.cc, so basically what this patch
> > does is
> > try to make the variables holding the array section base pointers to have
> > this
> > flag set.
> >
> > There is an another OMP_CLAUSE_MAP_POINTS_TO_READONLY set by front-ends on
> > the
> > associated pointer clauses if OMP_CLAUSE_MAP_READONLY is set.
> > Also a DECL_POINTS_TO_READONLY flag is set for VAR_DECLs when creating the
> > tmp
> > vars carrying these receiver references on the offloaded side. These
> > eventually get translated to SSA_NAME_POINTS_TO_READONLY_MEMORY.
>
>
> > This still doesn't always work as expected in terms of optimization:
> > struct pointer fields and Fortran arrays (kind of like C structs) which have
> > several accesses to create the pointer access on the receive/offloaded side,
> > and SRA appears to not work on these sequences, so gets in the way of much
> > redundancy elimination.
>
> I understand correctly that this is left as future work? Please add the test
> cases you have, XFAILed in some reasonable way.
>
>
> > Currently have one testcase where we can demonstrate 'readonly' can avoid
> > a clobber by function call.
>
> :-)
>
>
> > --- a/gcc/c/c-typeck.cc
> > +++ b/gcc/c/c-typeck.cc
> > @@ -14258,6 +14258,8 @@ handle_omp_array_sections (tree c, enum
> > c_omp_region_type ort)
> > OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
> > else
> > OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
> > + if (OMP_CLAUSE_MAP_READONLY (c))
> > + OMP_CLAUSE_MAP_POINTS_TO_READONLY (c2) = 1;
> > OMP_CLAUSE_MAP_IMPLICIT (c2) = OMP_CLAUSE_MAP_IMPLICIT (c);
> > if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
> > && !c_mark_addressable (t))
>
> > --- a/gcc/cp/semantics.cc
> > +++ b/gcc/cp/semantics.cc
> > @@ -5872,6 +5872,8 @@ handle_omp_array_sections (tree c, enum
> > c_omp_region_type ort)
> > }
> > else
> > OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
> > + if (OMP_CLAUSE_MAP_READONLY (c))
> > + OMP_CLAUSE_MAP_POINTS_TO_READONLY (c2) = 1;
> > OMP_CLAUSE_MAP_IMPLICIT (c2) = OMP_CLAUSE_MAP_IMPLICIT (c);
> > if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
> > && !cxx_mark_addressable (t))
>
> > --- a/gcc/fortran/trans-openmp.cc
> > +++ b/gcc/fortran/trans-openmp.cc
> > @@ -2524,6 +2524,8 @@ gfc_trans_omp_array_section (stmtblock_t *block,
> > gfc_exec_op op,
> > node3 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
> > OMP_CLAUSE_SET_MAP_KIND (node3, ptr_kind);
> > OMP_CLAUSE_DECL (node3) = gfc_conv_descriptor_data_get (decl);
> > + if (n->u.readonly)
> > + OMP_CLAUSE_MAP_POINTS_TO_READONLY (node3) = 1;
> > /* This purposely does not include GOMP_MAP_ALWAYS_POINTER. The
> > extra
> > cast prevents gimplify.cc from recognising it as being part of the
> > struct - and adding an 'alloc: for the 'desc.data' pointer, which
> > @@ -2559,6 +2561,8 @@ gfc_trans_omp_array_section (stmtblock_t *block,
> > gfc_exec_op op,
> > OMP_CLAUSE_MAP);
> > OMP_CLAUSE_SET_MAP_KIND (node3, ptr_kind);
> > OMP_CLAUSE_DECL (node3) = decl;
> > + if (n->u.readonly)
> > + OMP_CLAUSE_MAP_POINTS_TO_READONLY (node3) = 1;
> > }
>
> Could combine these two into one, after
> 'if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))' reconverges here, like
> where 'OMP_CLAUSE_SIZE (node3)' is set:
>
> > ptr2 = fold_convert (ptrdiff_type_node, ptr2);
> > OMP_CLAUSE_SIZE (node3) = fold_build2 (MINUS_EXPR, ptrdiff_type_node,
>
> Is 'n->u.readonly == OMP_CLAUSE_MAP_READONLY (node)'? If yes, would the
> latter be clearer to use as the 'if' expression (like in C, C++ front
> ends)?
>
> I see further additional 'OMP_CLAUSE_MAP' clauses synthesized, for
> example in 'gcc/cp/semantics.cc:handle_omp_array_sections', or
> 'gcc/fortran/trans-openmp.cc:gfc_trans_omp_array_section', also
> 'gcc/gimplify.cc'. I assume these are not relevant to have
> 'OMP_CLAUSE_MAP_READONLY' -> 'OMP_CLAUSE_MAP_POINTS_TO_READONLY'
> propagated? Actually, per your changes (see below), there is one
> 'OMP_CLAUSE_MAP_POINTS_TO_READONLY' propagation in
> 'gcc/gimplify.cc:build_omp_struct_comp_nodes'.
>
> Is the current situation re flag setting/propagation what was empirically
> necessary to make the test case work, or is it a systematic review? (The
> former is fine; I'd just like to know.)
>
> > --- a/gcc/gimple-expr.cc
> > +++ b/gcc/gimple-expr.cc
> > @@ -376,6 +376,8 @@ copy_var_decl (tree var, tree name, tree type)
> > DECL_CONTEXT (copy) = DECL_CONTEXT (var);
> > TREE_USED (copy) = 1;
> > DECL_SEEN_IN_BIND_EXPR_P (copy) = 1;
> > + if (VAR_P (var))
> > + DECL_POINTS_TO_READONLY (copy) = DECL_POINTS_TO_READONLY (var);
> > DECL_ATTRIBUTES (copy) = DECL_ATTRIBUTES (var);
> > if (DECL_USER_ALIGN (var))
> > {
>
> > --- a/gcc/gimplify.cc
> > +++ b/gcc/gimplify.cc
> > @@ -221,6 +221,7 @@ struct gimplify_omp_ctx
> > splay_tree variables;
> > hash_set<tree> *privatized_types;
> > tree clauses;
> > + hash_set<tree_operand_hash> *pt_readonly_ptrs;
> > /* Iteration variables in an OMP_FOR. */
> > vec<tree> loop_iter_var;
> > location_t location;
> > @@ -628,6 +629,15 @@ internal_get_tmp_var (tree val, gimple_seq *pre_p,
> > gimple_seq *post_p,
> > gimplify_expr (&val, pre_p, post_p, is_gimple_reg_rhs_or_call,
> > fb_rvalue);
> >
> > + bool pt_readonly = false;
> > + if (gimplify_omp_ctxp && gimplify_omp_ctxp->pt_readonly_ptrs)
> > + {
> > + tree ptr = val;
> > + if (TREE_CODE (ptr) == POINTER_PLUS_EXPR)
> > + ptr = TREE_OPERAND (ptr, 0);
> > + pt_readonly = gimplify_omp_ctxp->pt_readonly_ptrs->contains (ptr);
> > + }
>
> 'POINTER_PLUS_EXPR' is the only special thing we may run into, here?
> (Generally, I prefer 'if', 'else if, [...], 'else gcc_unreachable ()'.)
>
> > +
> > if (allow_ssa
> > && gimplify_ctxp->into_ssa
> > && is_gimple_reg_type (TREE_TYPE (val)))
> > @@ -639,9 +649,18 @@ internal_get_tmp_var (tree val, gimple_seq *pre_p,
> > gimple_seq *post_p,
> > if (name)
> > SET_SSA_NAME_VAR_OR_IDENTIFIER (t, create_tmp_var_name (name));
> > }
> > + if (pt_readonly)
> > + SSA_NAME_POINTS_TO_READONLY_MEMORY (t) = 1;
> > }
> > else
> > - t = lookup_tmp_var (val, is_formal, not_gimple_reg);
> > + {
> > + t = lookup_tmp_var (val, is_formal, not_gimple_reg);
> > + if (pt_readonly)
> > + {
> > + DECL_POINTS_TO_READONLY (t) = 1;
> > + gimplify_omp_ctxp->pt_readonly_ptrs->add (t);
> > + }
> > + }
> >
> > mod = build2 (INIT_EXPR, TREE_TYPE (t), t, unshare_expr (val));
> >
> > @@ -8906,6 +8925,8 @@ build_omp_struct_comp_nodes (enum tree_code code,
> > tree grp_start, tree grp_end,
> > OMP_CLAUSE_SET_MAP_KIND (c2, mkind);
> > OMP_CLAUSE_DECL (c2) = unshare_expr (OMP_CLAUSE_DECL (grp_end));
> > OMP_CLAUSE_CHAIN (c2) = NULL_TREE;
> > + if (OMP_CLAUSE_MAP_POINTS_TO_READONLY (grp_end))
> > + OMP_CLAUSE_MAP_POINTS_TO_READONLY (c2) = 1;
> > tree grp_mid = NULL_TREE;
> > if (OMP_CLAUSE_CHAIN (grp_start) != grp_end)
> > grp_mid = OMP_CLAUSE_CHAIN (grp_start);
>
> For my understanding, is this empirically necessary, or a systematic
> review?
>
> > @@ -11741,6 +11762,16 @@ gimplify_scan_omp_clauses (tree *list_p,
> > gimple_seq *pre_p,
> >
> > gimplify_omp_ctxp = outer_ctx;
> > }
> > + else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
> > + && (code == OACC_PARALLEL
> > + || code == OACC_KERNELS
> > + || code == OACC_SERIAL)
> > + && OMP_CLAUSE_MAP_POINTS_TO_READONLY (c))
> > + {
> > + if (ctx->pt_readonly_ptrs == NULL)
> > + ctx->pt_readonly_ptrs = new hash_set<tree_operand_hash> ();
> > + ctx->pt_readonly_ptrs->add (OMP_CLAUSE_DECL (c));
> > + }
> > if (notice_outer)
> > goto do_notice;
> > break;
>
> Also need to 'delete ctx->pt_readonly_ptrs;' somewhere.
>
> > --- a/gcc/omp-low.cc
> > +++ b/gcc/omp-low.cc
> > @@ -14098,6 +14098,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
> > omp_context *ctx)
> > if (ref_to_array)
> > x = fold_convert_loc (clause_loc, TREE_TYPE (new_var), x);
> > gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue);
> > + if (OMP_CLAUSE_MAP_POINTS_TO_READONLY (c) && VAR_P (x))
> > + DECL_POINTS_TO_READONLY (x) = 1;
> > if ((is_ref && !ref_to_array)
> > || ref_to_ptr)
> > {
>
> This is in the middle of the
> "Handle GOMP_MAP_FIRSTPRIVATE_{POINTER,REFERENCE} in second pass" code
> block. Again, for my understanding, is this empirically necessary, or a
> systematic review?
>
> > --- a/gcc/testsuite/c-c++-common/goacc/readonly-1.c
> > +++ b/gcc/testsuite/c-c++-common/goacc/readonly-1.c
> > @@ -19,8 +19,8 @@ int main (void)
> > 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 parallel
> > map\\(readonly,to:\\*s.ptr \\\[len: 64\\\]\\)
> > map\\(pt_readonly,attach_detach:s.ptr \\\[bias: 0\\\]\\)
> > map\\(readonly,to:x\\\[0\\\] \\\[len: 128\\\]\\)
> > map\\(pt_readonly,firstprivate:x \\\[pointer assign, bias: 0\\\]\\)" 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\\(pt_readonly,attach_detach:s.ptr \\\[bias: 0\\\]\\)
> > map\\(readonly,to:x\\\[0\\\] \\\[len: 128\\\]\\)
> > map\\(pt_readonly,firstprivate:x \\\[pointer assign, bias: 0\\\]\\)" 1
> > "original" { target { c++ } } } } */
> > /* { dg-final { scan-tree-dump-times "(?n)#pragma acc cache
> > \\(readonly:x\\\[0\\\] \\\[len: 128\\\]\\);$" 1 "original" } } */
>
> I suppose the new 'map(pt_readonly,attach_detach:s.ptr [bias: 0])' clause
> was previously "hidden" in '.+'? Please then change that in the first
> patch "[PATCH, OpenACC 2.7] readonly modifier support in front-ends", so
> that we can see here what actually is changing (only 'pt_readonly', I
> suppose).
>
> > --- /dev/null
> > +++ b/gcc/testsuite/c-c++-common/goacc/readonly-2.c
> > @@ -0,0 +1,15 @@
> > +/* { dg-additional-options "-O -fdump-tree-fre" } */
> > +
> > +#pragma acc routine
> > +extern void foo (int *ptr, int val);
> > +
> > +int main (void)
> > +{
> > + int r, a[32];
> > + #pragma acc parallel copyin(readonly: a[:32]) copyout(r)
> > + {
> > + foo (a, a[8]);
> > + r = a[8];
> > + }
> > +}
> > +/* { dg-final { scan-tree-dump-times "r\.\[_0-9\]+ =
> > MEM\\\[\[^_\]+_\[0-9\]+\\(ptro\\)\\\]\\\[8\\\];" 1 "fre1" } } */
>
> Please add a comment why 'fre1', and what generally is being checked
> here; that's not obvious to the casual reader. (That is, me in a few
> weeks.) ;-)
>
> Also add a scan for "before the optimization": two 'MEM's, I suppose?
>
> > --- a/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90
> > +++ b/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90
> > @@ -20,8 +20,8 @@ program main
> > !$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 parallel
> > map\\(readonly,to:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data
> > \\\[len: .+\\\]\\) map\\(pt_readonly,alloc:a.0 \\\[pointer assign, bias:
> > \\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) a.0\\\]\\)
> > map\\(readonly,to:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data
> > \\\[len: .+\\\]\\) map\\(pt_readonly,alloc:b \\\[pointer assign, bias:
> > \\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) b\\\]\\)"
> > 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\\(pt_readonly,alloc:a \\\[pointer assign, bias:
> > \\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) &a\\\]\\)
> > map\\(readonly,to:b\\\[\\(\\(integer\\(kind=8\\)\\) parm.*data -
> > \\(integer\\(kind=8\\)\\) &b\\) / 4\\\] \\\[len: .+\\\]\\)
> > map\\(pt_readonly,alloc:b \\\[pointer assign, bias:
> > \\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) &b\\\]\\)"
> > 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" } }
>
> Same comment as for 'c-c++-common/goacc/readonly-1.c'.
>
> > --- a/gcc/tree-pretty-print.cc
> > +++ b/gcc/tree-pretty-print.cc
> > @@ -907,6 +907,8 @@ dump_omp_clause (pretty_printer *pp, tree clause, int
> > spc, dump_flags_t flags)
> > pp_string (pp, "map(");
> > if (OMP_CLAUSE_MAP_READONLY (clause))
> > pp_string (pp, "readonly,");
> > + if (OMP_CLAUSE_MAP_POINTS_TO_READONLY (clause))
> > + pp_string (pp, "pt_readonly,");
> > switch (OMP_CLAUSE_MAP_KIND (clause))
> > {
> > case GOMP_MAP_ALLOC:
> > @@ -3436,6 +3438,8 @@ dump_generic_node (pretty_printer *pp, tree node, int
> > spc, dump_flags_t flags,
> > pp_string (pp, "(D)");
> > if (SSA_NAME_OCCURS_IN_ABNORMAL_PHI (node))
> > pp_string (pp, "(ab)");
> > + if (SSA_NAME_POINTS_TO_READONLY_MEMORY (node))
> > + pp_string (pp, "(ptro)");
> > break;
> >
> > case WITH_SIZE_EXPR:
>
> > --- a/gcc/tree-ssanames.cc
> > +++ b/gcc/tree-ssanames.cc
> > @@ -402,6 +402,9 @@ make_ssa_name_fn (struct function *fn, tree var, gimple
> > *stmt,
> > else
> > SSA_NAME_RANGE_INFO (t) = NULL;
> >
> > + if (VAR_P (var) && DECL_POINTS_TO_READONLY (var))
> > + SSA_NAME_POINTS_TO_READONLY_MEMORY (t) = 1;
> > +
> > SSA_NAME_IN_FREE_LIST (t) = 0;
> > SSA_NAME_IS_DEFAULT_DEF (t) = 0;
> > init_ssa_name_imm_use (t);
>
> > --- a/gcc/tree.h
> > +++ b/gcc/tree.h
> > @@ -1021,6 +1021,13 @@ extern void omp_clause_range_check_failed
> > (const_tree, const char *, int,
> > #define DECL_HIDDEN_STRING_LENGTH(NODE) \
> > (TREE_CHECK (NODE, PARM_DECL)->decl_common.decl_nonshareable_flag)
> >
> > +/* 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)
> > +
> > /* In a CALL_EXPR, means that the call is the jump from a thunk to the
> > thunked-to function. Be careful to avoid using this macro when one of
> > the
> > next two applies instead. */
> > @@ -1815,6 +1822,10 @@ class auto_suppress_location_wrappers
> > #define OMP_CLAUSE_MAP_READONLY(NODE) \
> > TREE_READONLY (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
> >
> > +/* Set if 'OMP_CLAUSE_DECL (NODE)' points to read-only memory. */
> > +#define OMP_CLAUSE_MAP_POINTS_TO_READONLY(NODE) \
> > + TREE_CONSTANT (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_))
>
> As in my "[PATCH, OpenACC 2.7] readonly modifier support in front-ends"
> review, please document how certain flags are used for OMP clauses.
>
>
> I note you're not actually using 'OMP_CLAUSE__CACHE__READONLY' anywhere
> -- but that's OK given the current 'gcc/gimplify.cc:gimplify_oacc_cache'.
> ;-)
>
>
> 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
