Re: [PATCH] OpenACC 2.6 manual deep copy support (attach/detach)

classic Classic list List threaded Threaded
5 messages Options
Reply | Threaded
Open this post in threaded view
|

Re: [PATCH] OpenACC 2.6 manual deep copy support (attach/detach)

Julian Brown-2
Hi Jakub,

Thank you for the review! Here's a new version of the patch.

On Fri, 7 Dec 2018 14:50:19 +0100
Jakub Jelinek <[hidden email]> wrote:

> On Fri, Nov 30, 2018 at 03:41:09AM -0800, Julian Brown wrote:
> > gcc/c-family/
> > * c-pragma.h (pragma_omp_clause): Add
> > PRAGMA_OACC_CLAUSE_ATTACH, PRAGMA_OACC_CLAUSE_DETACH.  
> ...
> > @@ -11804,9 +11808,12 @@ c_parser_omp_variable_list (c_parser
> > *parser, case OMP_CLAUSE_MAP:
> >      case OMP_CLAUSE_FROM:
> >      case OMP_CLAUSE_TO:
> > -      while (c_parser_next_token_is (parser, CPP_DOT))
> > +      while (c_parser_next_token_is (parser, CPP_DOT)
> > +     || c_parser_next_token_is (parser, CPP_DEREF))
> >   {
> >    location_t op_loc = c_parser_peek_token
> > (parser)->location;
> > +  if (c_parser_next_token_is (parser, CPP_DEREF))
> > +    t = build_simple_mem_ref (t);  
>
> This change is not ok, if OpenACC allows it in clauses, OpenMP 4.5
> does not and OpenMP 5.0 allows arbitrary lvalues that will need to be
> handled differently (still unimplemented).  So, this needs to be
> guarded for OpenACC only (perhaps for selected OpenACC clauses)?
I've added an "allow_deref" parameter to those functions. I guess that
can be extended/beautified when OpenMP 5.0 support comes along.

> > @@ -12632,6 +12631,8 @@ handle_omp_array_sections_1 (tree c, tree
> > t, vec<tree> &types, }
> >        t = TREE_OPERAND (t, 0);
> >      }
> > +  if (TREE_CODE (t) == MEM_REF)
> > +    t = TREE_OPERAND (t, 0);  
>
> Again, better guard this for OpenACC.  Maybe verify that
> mem_ref_offset is 0?

Done (though I'm not sure how to trigger that for testing!).

> > @@ -14163,6 +14214,8 @@ c_finish_omp_clauses (tree clauses, enum
> > c_omp_region_type ort) }
> >        if (remove)
> >   break;
> > +      if (TREE_CODE (t) == MEM_REF)
> > + t = TREE_OPERAND (t, 0);  
>
> Guard again?

OK.

> > @@ -31832,15 +31836,19 @@ cp_parser_omp_var_list_no_open (cp_parser
> > *parser, enum omp_clause_code kind, case OMP_CLAUSE_MAP:
> >      case OMP_CLAUSE_FROM:
> >      case OMP_CLAUSE_TO:
> > -      while (cp_lexer_next_token_is (parser->lexer,
> > CPP_DOT))
> > +      while (cp_lexer_next_token_is (parser->lexer,
> > CPP_DOT)
> > +     || cp_lexer_next_token_is (parser->lexer,
> > CPP_DEREF))  
>
> Ditto as for C.
Fixed, similarly.

> > @@ -4691,6 +4690,19 @@ handle_omp_array_sections_1 (tree c, tree t,
> > vec<tree> &types, if (low_bound == NULL_TREE)
> >      low_bound = integer_zero_node;
> >  
> > +  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
> > +      && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
> > +  || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
> > +    {
> > +      if (length != integer_one_node)
> > + {
> > +  error_at (OMP_CLAUSE_LOCATION (c),
> > +    OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
> > +    ? "array section in %<attach%> clause"
> > +    : "array section in %<detach%> clause");  
>
> So, are any array sections invalid, including e.g. [0:1] or say
> [5:] where size of the array is 6 elts, or what exactly is invalid?
IIUC, the restriction is that a single "attach" or "detach" clause may
only refer to a single pointer, so (possibly multiple dimensions of)
slices with single items are permitted. I've adjusted the error message
accordingly.

> > +      if (TREE_CODE (type) != POINTER_TYPE)
> > + {
> > +  error_at (OMP_CLAUSE_LOCATION (c),
> > +    OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
> > +    ? "expected pointer in %<attach%> clause"
> > +    : "expected pointer in %<detach%> clause");  
>
> Perhaps you can use %qs and omp_clause_name [OMP_CLAUSE_CODE (c)] ?

Fixed.

> >    }
> > + gfc_ref *array_ref = NULL;
> > + bool resolved = false;
> >   if (n->expr)
> >    {
> > -    if (!gfc_resolve_expr (n->expr)
> > +    array_ref = n->expr->ref;
> > +    resolved = gfc_resolve_expr (n->expr);
> > +
> > +    /* Look through component refs to find last
> > array
> > +       reference.  */
> > +    while (resolved
> > +   && array_ref
> > +   && (array_ref->type == REF_COMPONENT
> > +       || (array_ref->type == REF_ARRAY
> > +   && array_ref->next
> > +           && array_ref->next->type ==
> > REF_COMPONENT)))
> > +      array_ref = array_ref->next;  
>
> I'd guard this stuff for OpenACC only, keep what it did for OpenMP.
I've added a guard.

> > +  /* For OpenACC, pointers in structs should
> > trigger an
> > +     attach action.  */
> > +  if (ptr && (region_type & ORT_ACC) != 0)
> > +    {
> > +      /* Turning a GOMP_MAP_ALWAYS_POINTER clause
> > into a
> > + GOMP_MAP_ATTACH clause after we have
> > detected a case
> > + that needs a GOMP_MAP_STRUCT mapping
> > adding.  */
> > +      OMP_CLAUSE_SET_MAP_KIND (c,
> > + (code == OACC_EXIT_DATA) ? GOMP_MAP_DETACH
> > + :
> > GOMP_MAP_ATTACH);  
>
> Bad formatting, I'd suggest use a temporary with gomp_map_kind type.
Fixed.

> > +      has_attachments = true;
> > +    }
> >    if (n == NULL || (n->value & GOVD_MAP) == 0)
> >      {
> >        tree l = build_omp_clause
> > (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
> > -      OMP_CLAUSE_SET_MAP_KIND (l, GOMP_MAP_STRUCT);
> > +      OMP_CLAUSE_SET_MAP_KIND (l, attach
> > + ? GOMP_MAP_FORCE_PRESENT :
> > GOMP_MAP_STRUCT);  
>
> Likewise.
Fixed.

> >        if (!base_eq_orig_base)
> >   OMP_CLAUSE_DECL (l) = unshare_expr
> > (orig_base); else
> >   OMP_CLAUSE_DECL (l) = decl;
> > -      OMP_CLAUSE_SIZE (l) = size_int (1);
> > +      OMP_CLAUSE_SIZE (l) = attach
> > + ? (DECL_P (OMP_CLAUSE_DECL (l))
> > +     ? DECL_SIZE_UNIT (OMP_CLAUSE_DECL (l))
> > +     : TYPE_SIZE_UNIT (TREE_TYPE
> > (OMP_CLAUSE_DECL (l))))
> > + : size_int (1);  
>
> Again, bad formatting. = attach should be on a next line, the
> indentation is also weird, best like:
>      OMP_CLAUSE_SIZE (l)
> = (!attach
>   ? size_int (1)
>   : DECL_P (OMP_CLAUSE_DECL (l))
>   ? DECL_SIZE_UNIT (OMP_CLAUSE_DECL (l))
>   : TYPE_SIZE_UNIT (TREE_TYPE
> (OMP_CLAUSE_DECL (l))));
Fixed.

> > +/* { dg-final { scan-tree-dump-times "pragma omp target
> > oacc_enter_exit_data map.to:s .len: 32.." 1 "omplower" } } */ +/*
> > { dg-final { scan-tree-dump-times "pragma omp target oacc_data
> > map.tofrom:.z .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len:
> > 8.. map.tofrom:._1 .len: 40.. map.attach:s.a .len: 0.." 1
> > "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp
> > target oacc_parallel map.force_present:s .len: 32..
> > map.attach:s.e .len: 8.." 1 "omplower" } } */ +/* { dg-final
> > { scan-tree-dump-times "pragma omp target oacc_enter_exit_data
> > map.attach:a .len: 8.." 1 "omplower" } } */ +/* { dg-final
> > { scan-tree-dump-times "pragma omp target oacc_enter_exit_data
> > map.detach:a .len: 8.." 1 "omplower" } } */ +/* { dg-final
> > { scan-tree-dump-times "pragma omp target oacc_enter_exit_data
> > map.to:a .len: 8.." 1 "omplower" } } */ +/* { dg-final
> > { scan-tree-dump-times "pragma omp target oacc_enter_exit_data
> > map.force_present:s .len: 32.. map.detach:s.e .len: 8.." 1
> > "omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp
> > target oacc_data map.force_present:s .len: 32..
> > map.attach:s.e .len: 8.." 1 "omplower" } } */ +/* { dg-final
> > { scan-tree-dump-times "pragma omp target oacc_enter_exit_data
> > map.release:a .len: 8.." 1 "omplower" } } */ +/* { dg-final
> > { scan-tree-dump-times "pragma omp target oacc_enter_exit_data
> > finalize map.force_detach:a .len: 8.." 1 "omplower" } } */ +/*
> > { dg-final { scan-tree-dump-times "pragma omp target
> > oacc_enter_exit_data finalize map.force_present:s .len: 32..
> > map.force_detach:s.a .len: 8.." 1 "omplower" } } */  
>
> Aren't the lengths here heavily dependent on the target?  E.g. if it
> depends on sizeof (int) == 4, maybe the whole test needs to be
> guarded with { target int32 }
Fixed (by adding the { target int32 }).

> > @@ -918,8 +920,13 @@ struct splay_tree_key_s {
> >    uintptr_t tgt_offset;
> >    /* Reference count.  */
> >    uintptr_t refcount;
> > -  /* Dynamic reference count.  */
> > -  uintptr_t dynamic_refcount;
> > +  /* Reference counts beyond those that represent genuine
> > references in the
> > +     linked splay tree key/target memory structures, e.g. for
> > multiple OpenACC
> > +     "present increment" operations (via "acc enter data")
> > refering to the same
> > +     host-memory block.  */
> > +  uintptr_t virtual_refcount;
> > +  /* For a block with attached pointers, the attachment counters
> > for each.  */
> > +  unsigned short *attach_count;
> >    /* Pointer to the original mapping of "omp declare target link"
> > object.  */ splay_tree_key link_key;
> >  };  
>
> This is something I'm worried about a lot, the nodes keep growing way
> too much.  Is there a way to reuse some other field if it is of
> certain kind?
How about this -- it seems that the link_key is only used for OpenMP,
and the attach count is only needed for OpenACC. So the obvious thing
to do is probably to put those two together into a tagged union. The
question is where to put the tag?

Options are, I guess:

1. The high or low bits of the address.  Potentially non-portable, ugly.

2. Or, the virtual refcount is also only needed for OpenACC, so we can
   reserve a magic value for that field to act as a tag.

I've tried implementing the latter in the attached patch, and it seems
to work OK.

It occurs to me that "uintptr_t" is probably overkill for each of the
refcounts: we could probably reduce each of those to int32_t without
loss of too much generality, if we wanted to make further savings.

> Also, why unsigned short, can you only attach 65535 times?

A probably-misguided attempt to keep memory usage low (there's no such
restriction in the spec). I've switched this to uintptr_t, in line with
the current refcount sizes elsewhere.

Re-tested with offloading to nvptx. OK?

Thanks,

Julian

ChangeLog

        gcc/c-family/
        * c-pragma.h (pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_ATTACH,
        PRAGMA_OACC_CLAUSE_DETACH.

        gcc/c/
        * c-parser.c (c_parser_omp_clause_name): Add parsing of attach and
        detach clauses.
        (c_parser_omp_variable_list): Add ALLOW_DEREF parameter.  Allow deref
        (->) in variable lists if true.
        (c_parser_omp_var_list_parens): Add ALLOW_DEREF optional parameter.
        Pass to c_parser_omp_variable_list.
        (c_parser_oacc_data_clause): Support attach and detach clauses.  Update
        call to c_parser_omp_variable_list.
        (c_parser_omp_clause_reduction, c_parser_omp_clause_aligned)
        (c_parser_omp_clause_linear, c_parser_omp_clause_depend)
        (c_parser_omp_clause_map, c_parser_omp_clause_uniform): Update calls to
        c_parser_omp_variable_list.
        (c_parser_oacc_all_clauses): Support attach and detach clauses.
        (OACC_DATA_CLAUSE_MASK, OACC_ENTER_DATA_CLAUSE_MASK)
        (OACC_KERNELS_CLAUSE_MASK, OACC_PARALLEL_CLAUSE_MASK): Add
        PRAGMA_OACC_CLAUSE_ATTACH.
        (OACC_EXIT_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DETACH.
        * c-typeck.c (handle_omp_array_sections_1): Reject subarrays for attach
        and detach.  Support deref.
        (c_oacc_check_attachments): New function.
        (c_finish_omp_clauses): Check attach/detach arguments for being
        pointers using above.  Support deref.

        gcc/cp/
        * parser.c (cp_parser_omp_clause_name): Support attach and detach
        clauses.
        (cp_parser_omp_var_list_no_open): Add ALLOW_DEREF parameter.
        Parse deref if true.
        (cp_parser_omp_var_list): Add ALLOW_DEREF optional parameter.  Pass to
        cp_parser_omp_var_list_no_open.
        (cp_parser_oacc_data_clause): Support attach and detach clauses.  
        Update call to cp_parser_omp_var_list_no_open.
        (cp_parser_omp_clause_aligned, cp_parser_omp_clause_lastprivate)
        (cp_parser_omp_clause_linear, cp_parser_omp_clause_map): Update calls
        to cp_parser_omp_var_list_no_open.
        (cp_parser_oacc_all_clauses): Support attach and detach.
        (OACC_DATA_CLAUSE_MASK, OACC_ENTER_DATA_CLAUSE_MASK)
        (OACC_KERNELS_CLAUSE_MASK, OACC_PARALLEL_CLAUSE_MASK): Add
        PRAGMA_OACC_CLAUSE_ATTACH.
        (OACC_EXIT_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DETACH.
        * semantics.c (handle_omp_array_sections_1): Reject subarrays for
        attach and detach.
        (cp_oacc_check_attachments): New function.
        (finish_omp_clauses): Use above function.  Allow structure fields and
        class members to appear in OpenACC data clauses.  Support deref.

        gcc/fortran/
        * gfortran.h (gfc_omp_map_op): Add OMP_MAP_ATTACH, OMP_MAP_DETACH.
        * openmp.c (gfc_match_omp_variable_list): Add allow_derived parameter.
        Parse derived-type member accesses if true.
        (omp_mask2): Add OMP_CLAUSE_ATTACH, OMP_CLAUSE_DETACH.
        (gfc_match_omp_map_clause): Add allow_derived parameter.  Pass to
        gfc_match_omp_variable_list.
        (gfc_match_omp_clauses): Support attach and detach.  Support derived
        types for appropriate OpenACC directives.
        (OACC_PARALLEL_CLAUSES, OACC_KERNELS_CLAUSES, OACC_DATA_CLAUSES)
        (OACC_ENTER_DATA_CLAUSES): Add OMP_CLAUSE_ATTACH.
        (OACC_EXIT_DATA_CLAUSES): Add OMP_CLAUSE_DETACH.
        (check_symbol_not_pointer): Don't disallow pointer objects of derived
        type.
        (resolve_oacc_data_clauses): Don't disallow allocatable derived types.
        (resolve_omp_clauses): Perform duplicate checking only for non-derived
        type component accesses (plain variables and arrays or array sections).
        Support component refs.
        * trans-openmp.c (gfc_omp_privatize_by_reference): Support component
        refs.
        (gfc_trans_omp_clauses): Support component refs, attach and detach
        clauses.

        gcc/
        * gimplify.c (gimplify_omp_var_data): Add GOVD_MAP_HAS_ATTACHMENTS.
        (insert_struct_component_mapping): Support derived-type member mappings
        for arrays with descriptors which use GOMP_MAP_TO_PSET.
        (gimplify_scan_omp_clauses): Rewrite GOMP_MAP_ALWAYS_POINTER to
        GOMP_MAP_ATTACH for OpenACC struct/derived-type component pointers.
        Handle pointer mappings that use GOMP_MAP_TO_PSET.  Handle attach/detach
        clauses.
        (gimplify_adjust_omp_clauses_1): Skip adjustments for explicit
        attach/detach clauses.
        (gimplify_omp_target_update): Handle finalize for detach.
        * omp-low.c (lower_omp_target): Support GOMP_MAP_ATTACH,
        GOMP_MAP_DETACH, GOMP_MAP_FORCE_DETACH.
        * tree-pretty-print.c (dump_omp_clause): Likewise.

        gcc/include/
        * gomp-constants.h (GOMP_MAP_DEEP_COPY): Define.
        (gomp_map_kind): Add GOMP_MAP_ATTACH, GOMP_MAP_DETACH,
        GOMP_MAP_FORCE_DETACH.

        gcc/testsuite/
        * c-c++-common/goacc/mdc-1.c: New test.
        * c-c++-common/goacc/mdc-2.c: New test.
        * gcc.dg/goacc/mdc.C: New test.
        * gfortran.dg/goacc/data-clauses.f95: New test.
        * gfortran.dg/goacc/derived-types.f90: New test.
        * gfortran.dg/goacc/enter-exit-data.f95: New test.

        libgomp/
        * libgomp.h (struct target_var_desc): Add do_detach flag.
        (VREFCOUNT_LINK_KEY): New macro.
        (struct splay_tree_key_s): Put link_key and new attach_count field into
        a new union.  Substitute dynamic_refcount field for virtual_refcount.
        (struct acc_dispatch_t): Remove data_environ field.
        (enum gomp_map_vars_kind): Add GOMP_MAP_VARS_OPENACC_ENTER_DATA.
        (gomp_acc_insert_pointer): Remove prototype.
        (gomp_acc_remove_pointer): Update prototype.
        (struct gomp_coalesce_buf): Add forward declaration.
        (gomp_map_val, gomp_attach_pointer, gomp_detach_pointer): Add
        prototypes.
        * libgomp.map (OACC_2.6): New section. Add acc_attach, acc_attach_async,
        acc_detach, acc_detach_async, acc_detach_finalize,
        acc_detach_finalize_async.
        * oacc-async.c (goacc_remove_var_async): New function.
        * oacc-host.c (host_dispatch): Don't initialise removed data_environ
        field.
        * oacc-init.c (acc_shutdown_1): Use gomp_remove_var instead of
        gomp_unmap_vars to remove mappings by splay tree key instead of target
        memory descriptor.
        * oacc-int.h (splay_tree_key_s): Add forward declaration.
        (goacc_remove_var_async): Add prototype.
        * oacc-mem.c (lookup_dev_1): New function.
        (lookup_dev): Reimplement using above.
        (acc_free, acc_hostptr): Update calls to lookup_dev.
        (acc_map_data): Likewise.  Don't add to data_environ list.
        (acc_unmap_data): Remove call to gomp_unmap_vars.  Fix semantics to
        remove mapping, but not mapped data.
        (present_create_copy): Use virtual_refcount instead of
        dynamic_refcount.  Don't manipulate data_environ.  Fix target pointer
        return value.
        (delete_copyout): Update for virtual_refcount semantics.  Use
        goacc_remove_var_async for asynchronous delete/copyouts.
        (gomp_acc_insert_pointer): Remove function.
        (gomp_acc_remove_pointer): Reimplement.
        (acc_attach_async, acc_attach, goacc_detach_internal, acc_detach)
        (acc_detach_async, acc_detach_finalize, acc_detach_finalize_async): New
        functions.
        * oacc-parallel.c (find_pointer): Support attach/detach.  Make a little
        more strict.
        (GOACC_parallel_keyed): Use gomp_map_val to calculate device addresses.
        (GOACC_enter_exit_data): Support attach/detach and GOMP_MAP_STRUCT.
        Don't call gomp_acc_insert_pointer.
        * openacc.h (acc_attach, acc_attach_async, acc_detach)
        (acc_detach_async, acc_detach_finalize, acc_detach_finalize_async): Add
        prototypes.
        * target.c (limits.h): Include.
        (gomp_map_vars_existing): Initialise do_detach field of tgt_var_desc.
        (gomp_attach_pointer, gomp_detach_pointer): New functions.
        (gomp_map_val): Make global.
        (gomp_map_vars_async): Handle GOMP_MAP_VARS_OPENACC_ENTER_DATA.  Update
        for virtual_refcount semantics.  Support attach and detach.
        (gomp_remove_var): Free attach count array if present.
        (gomp_unmap_vars_async): Support detach and update for virtual_refcount
        semantics.  Disambiguate link_key/attach_count using virtual_refcount
        with magic value as a tag.
        (gomp_load_image_to_device): Zero-initialise virtual_refcount fields.
        (gomp_free_memmap): Remove function.
        (gomp_exit_data): Check virtual_refcount for tag value before using
        link_key.
        (omp_target_associate_ptr): Zero-initialise virtual_refcount and
        link_key splay tree key fields.
        (gomp_target_init): Don't initialise removed data_environ field.
        * testsuite/libgomp.oacc-c-c++-common/context-2.c: Use correct API to
        deallocate acc_copyin'd data.
        * testsuite/libgomp.oacc-c-c++-common/context-4.c: Likewise.
        * testsuite/libgomp.oacc-c-c++-common/deep-copy-1.c: New test.
        * testsuite/libgomp.oacc-c-c++-common/deep-copy-2.c: New test.
        * testsuite/libgomp.oacc-c-c++-common/deep-copy-3.c: New test.
        * testsuite/libgomp.oacc-c-c++-common/deep-copy-4.c: New test.
        * testsuite/libgomp.oacc-c-c++-common/deep-copy-5.c: New test.
        * testsuite/libgomp.oacc-c-c++-common/deep-copy-6.c: New test.
        * testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c: New test.
        * testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c: New test.
        * testsuite/libgomp.oacc-fortran/deep-copy-1.c: New test.
        * testsuite/libgomp.oacc-fortran/deep-copy-2.c: New test.
        * testsuite/libgomp.oacc-fortran/deep-copy-3.c: New test.
        * testsuite/libgomp.oacc-fortran/deep-copy-4.c: New test.
        * testsuite/libgomp.oacc-fortran/deep-copy-5.c: New test.
        * testsuite/libgomp.oacc-fortran/deep-copy-6.c: New test.
        * testsuite/libgomp.oacc-fortran/deep-copy-7.c: New test.
        * testsuite/libgomp.oacc-fortran/deep-copy-8.c: New test.
        * testsuite/libgomp.oacc-fortran/data-2.f90: Update test.
        * testsuite/libgomp.oacc-fortran/derived-type-1.f90: New test.
        * testsuite/libgomp.oacc-fortran/update-2.f90: New test.

attach-detach-3.diff (163K) Download Attachment
Reply | Threaded
Open this post in threaded view
|

Re: [PATCH] OpenACC 2.6 manual deep copy support (attach/detach)

Julian Brown-2
Hi!

Thanks again for review. Here's another iteration of the patch, with
your comments addressed.

One small thing, additionally: I've introduced a function
c_omp_map_clause_name in c-family/c-omp.c to return the name of OpenACC
map clauses -- since many of these use OMP_CLAUSE_MAP and then use
OMP_CLAUSE_MAP_KIND to distinguish between different mapping types,
just emitting 'map' to refer to several different clauses won't be very
meaningful to the user. I didn't see any existing way to get the right
names.

On Thu, 13 Dec 2018 11:57:05 +0100
Jakub Jelinek <[hidden email]> wrote:

> On Mon, Dec 10, 2018 at 07:41:37PM +0000, Julian Brown wrote:
> > @@ -11870,7 +11874,8 @@ c_parser_oacc_wait_list (c_parser *parser,
> > location_t clause_loc, tree list) static tree
> >  c_parser_omp_variable_list (c_parser *parser,
> >      location_t clause_loc,
> > -    enum omp_clause_code kind, tree list)
> > +    enum omp_clause_code kind, tree list,
> > +    bool allow_deref)  
>
> Make it bool allow_deref = false so that you don't have to change all
> callers?
Done.

> > @@ -12579,7 +12597,8 @@ c_parser_omp_clause_lastprivate (c_parser
> > *parser, tree list) }
> >   }
> >        tree nlist = c_parser_omp_variable_list (parser, loc,
> > -
> > OMP_CLAUSE_LASTPRIVATE, list);
> > +
> > OMP_CLAUSE_LASTPRIVATE, list,
> > +       false);
> >        c_parser_skip_until_found (parser, CPP_CLOSE_PAREN,
> > "expected %<)%>"); if (conditional)
> >   for (tree c = nlist; c != list; c = OMP_CLAUSE_CHAIN (c))  
>
> Like these etc.
Those bits reverted.

> > +  if (ort == C_ORT_ACC && TREE_CODE (t) == MEM_REF)
> > +    {
> > +      poly_int64 offset = mem_ref_offset (t).force_shwi ();
> > +      if (maybe_ne (offset, 0))  
>
> Just do if (maybe_ne (mem_ref_offset (t), 0)) ?

Fixed.

> > @@ -14432,6 +14491,16 @@ c_finish_omp_clauses (tree clauses, enum
> > c_omp_region_type ort) }
> >        if (remove)
> >   break;
> > +      if (ort == C_ORT_ACC && TREE_CODE (t) == MEM_REF)
> > +        {
> > +  poly_int64 offset = mem_ref_offset
> > (t).force_shwi ();
> > +  if (maybe_ne (offset, 0))  
>
> Likewise.
Fixed.

> > @@ -32111,7 +32115,7 @@ check_no_duplicate_clause (tree clauses,
> > enum omp_clause_code code,
> >  static tree
> >  cp_parser_omp_var_list_no_open (cp_parser *parser, enum
> > omp_clause_code kind,
> > - tree list, bool *colon)
> > + tree list, bool *colon, bool
> > allow_deref)  
>
> See above.
Fixed.

> > @@ -33560,7 +33579,7 @@ cp_parser_omp_clause_reduction (cp_parser
> > *parser, enum omp_clause_code kind, goto resync_fail;
> >  
> >    nlist = cp_parser_omp_var_list_no_open (parser, kind, list,
> > -  NULL);
> > +  NULL, false);
> >    for (c = nlist; c != list; c = OMP_CLAUSE_CHAIN (c))
> >      {
> >        OMP_CLAUSE_REDUCTION_CODE (c) = code;  
>
> See above.
Fixed.

> > --- a/libgomp/target.c
> > +++ b/libgomp/target.c
> > @@ -39,6 +39,9 @@
> >  #include <string.h>
> >  #include <assert.h>
> >  #include <errno.h>
> > +#ifdef RC_CHECKING
> > +#include <stdio.h>
> > +#endif  
>
> This doesn't belong here.
Oops! This leaked over from the refcount-checking patch via a bad
merge. Fixed.

> > @@ -1089,8 +1274,10 @@ gomp_remove_var (struct gomp_device_descr
> > *devicep, splay_tree_key k) {
> >    bool is_tgt_unmapped = false;
> >    splay_tree_remove (&devicep->mem_map, k);
> > -  if (k->link_key)
> > -    splay_tree_insert (&devicep->mem_map, (splay_tree_node)
> > k->link_key);
> > +  if (k->virtual_refcount == VREFCOUNT_LINK_KEY && k->u.link_key)
> > +    splay_tree_insert (&devicep->mem_map, (splay_tree_node)
> > k->u.link_key);
> > +  if (k->virtual_refcount != VREFCOUNT_LINK_KEY &&
> > k->u.attach_count)
> > +    free (k->u.attach_count);  
>
> So write
>   if (k->virtual_refcount == VREFCOUNT_LINK_KEY)
>     {
>       if (k->u.link_key)
> splay_tree_insert (&devicep->mem_map, (splay_tree_node)
> k->u.link_key); }
>   else if (k->u.attach_count)
>     free (k->u.attach_count);
> ?
I've done so.

Re-tested with offloading to nvptx. OK?

Thanks,

Julian

attach-detach-4.diff (161K) Download Attachment
Reply | Threaded
Open this post in threaded view
|

Re: [PATCH] OpenACC 2.6 manual deep copy support (attach/detach)

Thomas Schwinge-8
In reply to this post by Julian Brown-2
Hi!

While reviewing
<http://mid.mail-archive.com/20191003163505.49997-2-julian@...>
"OpenACC reference count overhaul", I've just now stumbled over one thing
that originally was designed here:

On 2018-12-10T19:41:37+0000, Julian Brown <[hidden email]> wrote:

> On Fri, 7 Dec 2018 14:50:19 +0100
> Jakub Jelinek <[hidden email]> wrote:
>
>> On Fri, Nov 30, 2018 at 03:41:09AM -0800, Julian Brown wrote:
>> > @@ -918,8 +920,13 @@ struct splay_tree_key_s {
>> >    uintptr_t tgt_offset;
>> >    /* Reference count.  */
>> >    uintptr_t refcount;
>> > -  /* Dynamic reference count.  */
>> > -  uintptr_t dynamic_refcount;
>> > +  /* Reference counts beyond those that represent genuine references in the
>> > +     linked splay tree key/target memory structures, e.g. for multiple OpenACC
>> > +     "present increment" operations (via "acc enter data") refering to the same
>> > +     host-memory block.  */
>> > +  uintptr_t virtual_refcount;
>> > +  /* For a block with attached pointers, the attachment counters for each.  */
>> > +  unsigned short *attach_count;
>> >    /* Pointer to the original mapping of "omp declare target link" object.  */
>> >    splay_tree_key link_key;
>> >  };  
>>
>> This is something I'm worried about a lot, the nodes keep growing way
>> too much.
Is that just a would-be-nice-to-avoid, or is it an actual problem?

If the latter, can we maybe move some data into on-the-side data
structures, say an associative array keyed by [something suitable]?  I
would assume that compared to actual host to/from device data movement
(or even lookup etc.), lookup of values from such an associative array
should be relatively cheap?

I'm bringing this up, because:

>> Is there a way to reuse some other field if it is of
>> certain kind?
>
> How about this -- it seems that the link_key is only used for OpenMP,

So, is that actually correct?  Per my understanding, for the OpenACC
'link' clause we uses 'GOMP_MAP_LINK', which sets "omp declare target
link", and thus:

> and the attach count is only needed for OpenACC. So the obvious thing
> to do is probably to put those two together into a tagged union. The
> question is where to put the tag?
>
> Options are, I guess:
>
> 1. The high or low bits of the address.  Potentially non-portable, ugly.
>
> 2. Or, the virtual refcount is also only needed for OpenACC, so we can
>    reserve a magic value for that field to act as a tag.
>
> I've tried implementing the latter in the attached patch, and it seems
> to work OK.
... this is not actually feasible?

It's certainly possible that we're totally lacking sufficient testsuite
coverage, and that there are issues in the 'link' implementation
(<https://gcc.gnu.org/PR81689> "libgomp.c/target-link-1.c fails for
nvptx: #pragma omp target link not implemented" comes to mind
immediatelly, and certainly for OpenACC I used to be aware of additional
issues; I think I intended to use that mechanism for Fortran
'allocatable' with OpenACC 'declare'), but the libgomp handling to me
seems reasonable upon quick review -- just that we need to keep it alive
for OpenACC, too, unless I'm confused?


Simplifying the libgomp code to avoid the 'VREFCOUNT_LINK_KEY' toggle
flag, and not putting 'link_key' into an union together with
'attach_count', that should -- I hope -- resolve/obsolete some of the
questions raised in my late-night pre-Christmas 2018 review,
<https://gcc.gnu.org/ml/gcc-patches/2018-12/msg01620.html>, where I'm now
not sure yet whether all my questions have been addressed (or disputed,
but I didn't hear anything) in the recent -- split-out, thanks! --
version of this patch,
<http://mid.mail-archive.com/20191003163505.49997-2-julian@...>
"OpenACC reference count overhaul".


Grüße
 Thomas

signature.asc (671 bytes) Download Attachment
Reply | Threaded
Open this post in threaded view
|

Re: [PATCH] OpenACC 2.6 manual deep copy support (attach/detach)

Julian Brown-2
Hi,

This version of the patch has been adjusted for changes to the
reference-count overhaul patch last posted here:

https://gcc.gnu.org/ml/gcc-patches/2019-10/msg02041.html

In particular, it incorporates the idea alluded to in that patch
relating to adding the new "attach_count" field to the memory-mapping
splay tree key type without growing that structure in the common case
(i.e. when structure components are not being mapped, or for OpenMP).
In short, a new auxiliary structure is added containing the previous
"link_key" and "attach_count" fields: so, you can either have both
pointers (though of course one of them may be NULL), or in the common
case no aux pointer at all, so no growth in the base struct size.

I've also addressed a couple of Thomas's earlier review comments from
December 2018:

https://gcc.gnu.org/ml/gcc-patches/2018-12/msg01620.html

And also, as a minor point, the patch reincorporates the patch adding
c_omp_map_clause_name, since it refers to the newly-added
GOMP_MAP_ATTACH values so the two patches are mutually dependent. That
one was posted here:

https://gcc.gnu.org/ml/gcc-patches/2019-10/msg00443.html

On Fri, 18 Oct 2019 18:47:08 +0200
Thomas Schwinge <[hidden email]> wrote:

> Hi!
>
> While reviewing
> <http://mid.mail-archive.com/20191003163505.49997-2-julian@...>
> "OpenACC reference count overhaul", I've just now stumbled over one
> thing that originally was designed here:
>
> On 2018-12-10T19:41:37+0000, Julian Brown <[hidden email]>
> wrote:
> > On Fri, 7 Dec 2018 14:50:19 +0100
> > Jakub Jelinek <[hidden email]> wrote:
> >  
> >> On Fri, Nov 30, 2018 at 03:41:09AM -0800, Julian Brown wrote:  
> >> > @@ -918,8 +920,13 @@ struct splay_tree_key_s {
> >> >    uintptr_t tgt_offset;
> >> >    /* Reference count.  */
> >> >    uintptr_t refcount;
> >> > -  /* Dynamic reference count.  */
> >> > -  uintptr_t dynamic_refcount;
> >> > +  /* Reference counts beyond those that represent genuine
> >> > references in the
> >> > +     linked splay tree key/target memory structures, e.g. for
> >> > multiple OpenACC
> >> > +     "present increment" operations (via "acc enter data")
> >> > refering to the same
> >> > +     host-memory block.  */
> >> > +  uintptr_t virtual_refcount;
> >> > +  /* For a block with attached pointers, the attachment
> >> > counters for each.  */
> >> > +  unsigned short *attach_count;
> >> >    /* Pointer to the original mapping of "omp declare target
> >> > link" object.  */ splay_tree_key link_key;
> >> >  };    
> >>
> >> This is something I'm worried about a lot, the nodes keep growing
> >> way too much.  
>
> Is that just a would-be-nice-to-avoid, or is it an actual problem?
>
> If the latter, can we maybe move some data into on-the-side data
> structures, say an associative array keyed by [something suitable]?  I
> would assume that compared to actual host to/from device data movement
> (or even lookup etc.), lookup of values from such an associative array
> should be relatively cheap?
I'd be extremely wary of adding a completely separate off-the-side
structure to keep track of attachment counters: the reference-counting
behaviour is already complicated enough, and the risk of messing things
up with another indirectly-linked structure to keep track of is too
high (never mind the extra runtime overhead). With the approach in this
patch, at least the extra info for link_key/attach_count is directly
accessible from the splay tree key struct via pointer indirection.

This version entails slight additional overhead (another malloc'd
block and another pointer indirection) for the link_key field (and also
for the attach_count pointer). I've not benchmarked memory use or
performance though, so I'm not sure how much impact this has on real
code.

> I'm bringing this up, because:
>
> >> Is there a way to reuse some other field if it is of
> >> certain kind?  
> >
> > How about this -- it seems that the link_key is only used for
> > OpenMP,  
>
> So, is that actually correct?  Per my understanding, for the OpenACC
> 'link' clause we uses 'GOMP_MAP_LINK', which sets "omp declare target
> link", and thus:
>
> > and the attach count is only needed for OpenACC. So the obvious
> > thing to do is probably to put those two together into a tagged
> > union. The question is where to put the tag?
> >
> > Options are, I guess:
> >
> > 1. The high or low bits of the address.  Potentially non-portable,
> > ugly.
> >
> > 2. Or, the virtual refcount is also only needed for OpenACC, so we
> > can reserve a magic value for that field to act as a tag.
> >
> > I've tried implementing the latter in the attached patch, and it
> > seems to work OK.  
>
> ... this is not actually feasible?
I'd missed that link_key could be used for OpenACC as well as OpenMP
(obviously). This version of the patch avoids that problem.

Is this version OK? Re-tested with offloading to nvptx.

Thanks,

Julian

ChangeLog

2019-11-06  Julian Brown  <[hidden email]>
            Thomas Schwinge  <[hidden email]>

    gcc/c-family/
    * c-common.h (c_omp_map_clause_name): Add prototype.
    * c-omp.c (c_omp_map_clause_name): New function.
    * c-pragma.h (pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_ATTACH and
    PRAGMA_OACC_CLAUSE_DETACH.

    gcc/c/
    * c-parser.c (c_parser_omp_clause_name): Add parsing of attach and
    detach clauses.
    (c_parser_omp_variable_list): Add ALLOW_DEREF optional parameter.
    Allow deref (->) in variable lists if true.
    (c_parser_omp_var_list_parens): Add ALLOW_DEREF optional parameter.
    Pass to c_parser_omp_variable_list.
    (c_parser_oacc_data_clause): Support attach and detach clauses.  Update
    call to c_parser_omp_variable_list.
    (c_parser_oacc_all_clauses): Support attach and detach clauses.
    (OACC_DATA_CLAUSE_MASK, OACC_ENTER_DATA_CLAUSE_MASK)
    (OACC_KERNELS_CLAUSE_MASK, OACC_PARALLEL_CLAUSE_MASK): Add
    PRAGMA_OACC_CLAUSE_ATTACH.
    (OACC_EXIT_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DETACH.
    * c-typeck.c (handle_omp_array_sections_1): Reject subarrays for attach
    and detach.  Support deref.
    (handle_omp_array_sections): Use GOMP_MAP_ATTACH_DETACH instead of
    GOMP_MAP_ALWAYS_POINTER for OpenACC.
    (c_oacc_check_attachments): New function.
    (c_finish_omp_clauses): Check attach/detach arguments for being
    pointers using above.  Support deref.

    gcc/cp/
    * parser.c (cp_parser_omp_clause_name): Support attach and detach
    clauses.
    (cp_parser_omp_var_list_no_open): Add ALLOW_DEREF optional parameter.
    Parse deref if true.
    (cp_parser_omp_var_list): Add ALLOW_DEREF optional parameter.  Pass to
    cp_parser_omp_var_list_no_open.
    (cp_parser_oacc_data_clause): Support attach and detach clauses.
    Update call to cp_parser_omp_var_list_no_open.
    (cp_parser_oacc_all_clauses): Support attach and detach.
    (OACC_DATA_CLAUSE_MASK, OACC_ENTER_DATA_CLAUSE_MASK)
    (OACC_KERNELS_CLAUSE_MASK, OACC_PARALLEL_CLAUSE_MASK): Add
    PRAGMA_OACC_CLAUSE_ATTACH.
    (OACC_EXIT_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DETACH.
    * semantics.c (handle_omp_array_sections_1): Reject subarrays for
    attach and detach.
    (handle_omp_array_sections): Use GOMP_MAP_ATTACH_DETACH instead of
    GOMP_MAP_ALWAYS_POINTER for OpenACC.
    (cp_oacc_check_attachments): New function.
    (finish_omp_clauses): Use above function.  Allow structure fields and
    class members to appear in OpenACC data clauses.  Support
    GOMP_MAP_ATTACH_DETACH.  Support deref.

    gcc/fortran/
    * gfortran.h (gfc_omp_map_op): Add OMP_MAP_ATTACH, OMP_MAP_DETACH.
    * openmp.c (gfc_match_omp_variable_list): Add allow_derived parameter.
    Parse derived-type member accesses if true.
    (omp_mask2): Add OMP_CLAUSE_ATTACH and OMP_CLAUSE_DETACH.
    (gfc_match_omp_map_clause): Add allow_derived parameter.  Pass to
    gfc_match_omp_variable_list.
    (gfc_match_omp_clauses): Support attach and detach.  Support derived
    types for appropriate OpenACC directives.
    (OACC_PARALLEL_CLAUSES, OACC_KERNELS_CLAUSES, OACC_DATA_CLAUSES)
    (OACC_ENTER_DATA_CLAUSES): Add OMP_CLAUSE_ATTACH.
    (OACC_EXIT_DATA_CLAUSES): Add OMP_CLAUSE_DETACH.
    (check_symbol_not_pointer): Don't disallow pointer objects of derived
    type.
    (resolve_oacc_data_clauses): Don't disallow allocatable derived types.
    (resolve_omp_clauses): Perform duplicate checking only for non-derived
    type component accesses (plain variables and arrays or array sections).
    Support component refs.
    * trans-expr.c (gfc_conv_component_ref,
    conv_parent_component_references): Make global.
    (gfc_auto_dereference_var): New function, broken out of...
    (gfc_conv_variable): ...here.  Call above function.
    * trans-openmp.c (gfc_omp_privatize_by_reference): Support component
    refs.
    (gfc_trans_omp_array_section): New function, broken out of...
    (gfc_trans_omp_clauses): ...here.  Support component refs/derived
    types, attach and detach clauses.
    * trans.h (gfc_conv_component_ref, conv_parent_component_references,
    gfc_auto_dereference_var): Add prototypes.

    gcc/
    * gimplify.c (gimplify_omp_var_data): Add GOVD_MAP_HAS_ATTACHMENTS.
    (insert_struct_comp_map): Support derived-type member mappings
    for arrays with descriptors which use GOMP_MAP_TO_PSET.  Support
    GOMP_MAP_ATTACH_DETACH.
    (gimplify_scan_omp_clauses): Tidy up OACC_ENTER_DATA/OACC_EXIT_DATA
    mappings.  Handle attach/detach clauses and component references.
    (gimplify_adjust_omp_clauses_1): Skip adjustments for explicit
    attach/detach clauses.
    (gimplify_omp_target_update): Handle finalize for detach.
    * omp-low.c (lower_omp_target): Support GOMP_MAP_ATTACH,
    GOMP_MAP_DETACH, GOMP_MAP_FORCE_DETACH.
    * tree-pretty-print.c (dump_omp_clause): Likewise, plus
    GOMP_MAP_ATTACH_DETACH.

    include/
    * gomp-constants.h (GOMP_MAP_FLAG_SPECIAL_4, GOMP_MAP_DEEP_COPY):
    Define.
    (gomp_map_kind): Add GOMP_MAP_ATTACH, GOMP_MAP_DETACH,
    GOMP_MAP_FORCE_DETACH.

    gcc/testsuite/
    * c-c++-common/goacc/deep-copy-arrayofstruct.c: New test.
    * c-c++-common/goacc/mdc-1.c: New test.
    * c-c++-common/goacc/mdc-2.c: New test.
    * gcc.dg/goacc/mdc.C: New test.
    * gfortran.dg/goacc/derived-types.f90: New test.
    * gfortran.dg/goacc/derived-types-2.f90: New test.
    * gfortran.dg/goacc/data-clauses.f95: Adjust for expected errors.
    * gfortran.dg/goacc/enter-exit-data.f95: Likewise.

    libgomp/
    * libgomp.h (struct target_var_desc): Add do_detach flag.
    (struct splay_tree_aux): New.
    (struct splay_tree_key_s): Replace link_key field with aux pointer.
    (gomp_attach_pointer, gomp_detach_pointer): Add prototypes.
    * libgomp.map (OACC_2.6): New section. Add acc_attach,
    acc_attach_async, acc_detach, acc_detach_async, acc_detach_finalize,
    acc_detach_finalize_async.
    * oacc-init.c (acc_shutdown_1): Free aux block if present.
    * oacc-mem.c (acc_attach_async, acc_attach, goacc_detach_internal,
    acc_detach, acc_detach_async, acc_detach_finalize,
    acc_detach_finalize_async): New functions.
    * oacc-parallel.c (find_group_last): Add SIZES parameter. Support
    struct components.  Tidy up and add some new checks.
    (goacc_enter_data_internal): Update call to find_group_last.
    (goacc_exit_data_internal): Support detach operations and
    GOMP_MAP_STRUCT.
    (GOACC_enter_exit_data): Handle initial GOMP_MAP_STRUCT or
    GOMP_MAP_FORCE_PRESENT in finalization detection code.  Handle
    attach/detach in enter/exit data detection code.
    * openacc.h (acc_attach, acc_attach_async, acc_detach,
    (acc_detach_async, acc_detach_finalize, acc_detach_finalize_async): Add
    prototypes.
    * target.c (dump_tgt): Support aux field.
    (gomp_map_vars_existing): Initialise do_detach field of tgt_var_desc.
    (gomp_attach_pointer, gomp_detach_pointer): New functions.
    (gomp_map_vars_internal): Support attach and detach.
    (gomp_remove_var_internal): Free aux block and attachment counts if
    present.
    (gomp_unmap_vars_internal): Support detach.
    (gomp_load_image_to_device): Zero-initialise aux field instead of
    link_key field.
    (gomp_exit_data): Handle link key in aux field.  Free aux field when
    appropriate.
    (omp_target_associate_ptr): Zero-initialize aux field instead of
    link_key.
    * testsuite/libgomp.oacc-c-c++-common/deep-copy-1.c: New test.
    * testsuite/libgomp.oacc-c-c++-common/deep-copy-2.c: New test.
    * testsuite/libgomp.oacc-c-c++-common/deep-copy-3.c: New test.
    * testsuite/libgomp.oacc-c-c++-common/deep-copy-4.c: New test.
    * testsuite/libgomp.oacc-c-c++-common/deep-copy-5.c: New test.
    * testsuite/libgomp.oacc-c-c++-common/deep-copy-6.c: New test.
    * testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c: New test.
    * testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c: New test.
    * testsuite/libgomp.oacc-c-c++-common/deep-copy-9.c: New test.
    * testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c: New test.
    * testsuite/libgomp.oacc-c-c++-common/deep-copy-11.c: New test.
    * testsuite/libgomp.oacc-c-c++-common/deep-copy-14.c: New test.
    * testsuite/libgomp.oacc-c++/deep-copy-12.C: New test.
    * testsuite/libgomp.oacc-c++/deep-copy-13.C: New test.
    * testsuite/libgomp.oacc-fortran/deep-copy-1.f90: New test.
    * testsuite/libgomp.oacc-fortran/deep-copy-2.f90: New test.
    * testsuite/libgomp.oacc-fortran/deep-copy-3.f90: New test.
    * testsuite/libgomp.oacc-fortran/deep-copy-4.f90: New test.
    * testsuite/libgomp.oacc-fortran/deep-copy-5.f90: New test.
    * testsuite/libgomp.oacc-fortran/deep-copy-6.f90: New test.
    * testsuite/libgomp.oacc-fortran/deep-copy-7.f90: New test.
    * testsuite/libgomp.oacc-fortran/deep-copy-8.f90: New test.
    * testsuite/libgomp.oacc-fortran/derived-type-1.f90: New test.
    * testsuite/libgomp.oacc-fortran/derivedtype-1.f95: New test.
    * testsuite/libgomp.oacc-fortran/derivedtype-2.f95: New test.
    * testsuite/libgomp.oacc-fortran/multidim-slice.f95: New test.
    * testsuite/libgomp.oacc-fortran/update-2.f90: New test.

attach-detach-fsf-20191106.diff (177K) Download Attachment
Reply | Threaded
Open this post in threaded view
|

Re: [PATCH] OpenACC 2.6 manual deep copy support (attach/detach)

Tobias Burnus-3
Hi Julian,

On 11/23/19 12:42 AM, Julian Brown wrote:
>
>              gcc/fortran/
>              * trans-expr.c
>              (gfc_auto_dereference_var): New function, broken out of...
> […]
> +gfc_auto_dereference_var (location_t loc, gfc_symbol *sym, tree var,
> +  bool descriptor_only_p, bool is_classarray)
> […]
> + var = build_fold_indirect_ref_loc (input_location, var);
As your patch takes a location as argument – it also makes sense to use
that location. (Alternatively, one could remove the argument as both
callers explicitly pass 'input_location' as argument.)

One could do either way, but the current variant does not make sense –
and, additionally, the current variant causes a compile-time warning.

See also my OG9 commit 500483e6ced44e2e0fea6a37e4f8c267ebaf826a where do
s/input_location/loc/g in that function.

Cheers,

Tobias