Re: [PATCH, OpenACC 2.7] Connect readonly modifier to points-to analysis

[Thomas Schwinge <thomas@codesourcery.com>]

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-b30adbd5901a@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:

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